





<!DOCTYPE html>
<html class="writer-html5" lang="zh-CN" >
<head>
  <meta charset="utf-8">
  
  <meta name="viewport" content="width=device-width, initial-scale=1.0">
  
  <title>Tuning High Performance Convolution on NVIDIA GPUs &mdash; tvm 0.8.dev1982 文档</title>
  

  
  <link rel="stylesheet" href="https://maxcdn.bootstrapcdn.com/bootstrap/4.0.0/css/bootstrap.min.css" integrity="sha384-Gn5384xqQ1aoWXA+058RXPxPg6fy4IWvTNh0E263XmFcJlSAwiGgFAW/dAiS6JXm" crossorigin="anonymous">
  <link rel="stylesheet" href="../../_static/css/theme.css" type="text/css" />
  <link rel="stylesheet" href="../../_static/pygments.css" type="text/css" />
  <link rel="stylesheet" href="../../_static/css/theme.css" type="text/css" />
  <link rel="stylesheet" href="../../_static/gallery.css" type="text/css" />
  <link rel="stylesheet" href="../../_static/pygments.css" type="text/css" />
  <link rel="stylesheet" href="../../_static/css/tlcpack_theme.css" type="text/css" />

  
  
    <link rel="shortcut icon" href="../../_static/tvm-logo-square.png"/>
  

  
  
  
  
    
      <script type="text/javascript" id="documentation_options" data-url_root="../../" src="../../_static/documentation_options.js"></script>
        <script data-url_root="../../" id="documentation_options" src="../../_static/documentation_options.js"></script>
        <script src="../../_static/jquery.js"></script>
        <script src="../../_static/underscore.js"></script>
        <script src="../../_static/doctools.js"></script>
        <script src="../../_static/translations.js"></script>
    
    <script type="text/javascript" src="../../_static/js/theme.js"></script>

    
    <script type="text/javascript" src="../../_static/js/tlcpack_theme.js"></script>
    <link rel="index" title="索引" href="../../genindex.html" />
    <link rel="search" title="搜索" href="../../search.html" />
    <link rel="next" title="Auto-tuning a Convolutional Network for NVIDIA GPU" href="tune_relay_cuda.html" />
    <link rel="prev" title="Auto-Tune with Templates and AutoTVM" href="index.html" /> 
</head>

<body class="wy-body-for-nav">

   
  <div class="wy-grid-for-nav">
    
    
<header class="header">
    <div class="innercontainer">
      <div class="headerInner d-flex justify-content-between align-items-center">
          <div class="headerLogo">
               <a href="https://tvm.apache.org/"><img src=https://tvm.apache.org/assets/images/logo.svg alt="logo"></a>
          </div>

          <div id="headMenu" class="headerNav">
            <button type="button" id="closeHeadMenu" class="navCloseBtn"><img src="../../_static/img/close-icon.svg" alt="Close"></button>
             <ul class="nav">
                <li class="nav-item">
                   <a class="nav-link" href=https://tvm.apache.org/community>Community</a>
                </li>
                <li class="nav-item">
                   <a class="nav-link" href=https://tvm.apache.org/download>Download</a>
                </li>
                <li class="nav-item">
                   <a class="nav-link" href=https://tvm.apache.org/vta>VTA</a>
                </li>
                <li class="nav-item">
                   <a class="nav-link" href=https://tvm.apache.org/blog>Blog</a>
                </li>
                <li class="nav-item">
                   <a class="nav-link" href=https://tvm.apache.org/docs>Docs</a>
                </li>
                <li class="nav-item">
                   <a class="nav-link" href=https://tvmconf.org>Conference</a>
                </li>
                <li class="nav-item">
                   <a class="nav-link" href=https://github.com/apache/tvm/>Github</a>
                </li>
                <li class="nav-item">
                   <a class="nav-link" href=https://tvmchinese.github.io/declaration_zh_CN.html>About-Translators</a>
                </li>
             </ul>
               <div class="responsivetlcdropdown">
                 <button type="button" class="btn-link">
                   ASF
                 </button>
                 <ul>
                     <li>
                       <a href=https://apache.org/>Apache Homepage</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/licenses/>License</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/foundation/sponsorship.html>Sponsorship</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/security/>Security</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/foundation/thanks.html>Thanks</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/events/current-event>Events</a>
                     </li>
                     <li>
                       <a href=https://www.zhihu.com/column/c_1429578595417563136>Zhihu</a>
                     </li>
                 </ul>
               </div>
          </div>
            <div class="responsiveMenuIcon">
              <button type="button" id="menuBtn" class="btn-menu"><img src="../../_static/img/menu-icon.svg" alt="Menu Icon"></button>
            </div>

            <div class="tlcDropdown">
              <div class="dropdown">
                <button type="button" class="btn-link dropdown-toggle" data-toggle="dropdown" aria-haspopup="true" aria-expanded="false">
                  ASF
                </button>
                <div class="dropdown-menu dropdown-menu-right">
                  <ul>
                     <li>
                       <a href=https://apache.org/>Apache Homepage</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/licenses/>License</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/foundation/sponsorship.html>Sponsorship</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/security/>Security</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/foundation/thanks.html>Thanks</a>
                     </li>
                     <li>
                       <a href=https://www.apache.org/events/current-event>Events</a>
                     </li>
                     <li>
                       <a href=https://www.zhihu.com/column/c_1429578595417563136>Zhihu</a>
                     </li>
                  </ul>
                </div>
              </div>
          </div>
       </div>
    </div>
 </header>
 
    <nav data-toggle="wy-nav-shift" class="wy-nav-side fixed">
      <div class="wy-side-scroll">
        <div class="wy-side-nav-search" >
          

          
            <a href="../../index.html">
          

          
            
            <img src="../../_static/tvm-logo-small.png" class="logo" alt="Logo"/>
          
          </a>

          
            
            
                <div class="version">
                  0.8.dev1982
                </div>
            
          

          
<div role="search">
  <form id="rtd-search-form" class="wy-form" action="../../search.html" method="get">
    <input type="text" name="q" placeholder="Search docs" />
    <input type="hidden" name="check_keywords" value="yes" />
    <input type="hidden" name="area" value="default" />
  </form>
</div>

          
        </div>

        
        <div class="wy-menu wy-menu-vertical" data-spy="affix" role="navigation" aria-label="main navigation">
          
            
            
              
            
            
              <p class="caption" role="heading"><span class="caption-text">如何开始</span></p>
<ul>
<li class="toctree-l1"><a class="reference internal" href="../../install/index.html">安装 TVM</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../contribute/index.html">贡献者指南</a></li>
</ul>
<p class="caption" role="heading"><span class="caption-text">用户引导</span></p>
<ul class="current">
<li class="toctree-l1"><a class="reference internal" href="../../tutorial/index.html">User Tutorial</a></li>
<li class="toctree-l1 current"><a class="reference internal" href="../index.html">How To Guides</a><ul class="current">
<li class="toctree-l2"><a class="reference internal" href="../compile_models/index.html">编译深度学习模型</a></li>
<li class="toctree-l2"><a class="reference internal" href="../deploy/index.html">TVM 部署模型和集成</a></li>
<li class="toctree-l2"><a class="reference internal" href="../work_with_relay/index.html">Work With Relay</a></li>
<li class="toctree-l2"><a class="reference internal" href="../work_with_schedules/index.html">Work With Tensor Expression and Schedules</a></li>
<li class="toctree-l2"><a class="reference internal" href="../optimize_operators/index.html">优化张量算子</a></li>
<li class="toctree-l2 current"><a class="reference internal" href="index.html">Auto-Tune with Templates and AutoTVM</a><ul class="current">
<li class="toctree-l3 current"><a class="current reference internal" href="#">Tuning High Performance Convolution on NVIDIA GPUs</a><ul>
<li class="toctree-l4"><a class="reference internal" href="#install-dependencies">安装依赖</a></li>
<li class="toctree-l4"><a class="reference internal" href="#step-1-define-the-search-space">Step 1:  Define the search space</a></li>
<li class="toctree-l4"><a class="reference internal" href="#step-2-search-through-the-space">Step 2:  Search through the space</a></li>
</ul>
</li>
<li class="toctree-l3"><a class="reference internal" href="tune_relay_cuda.html">Auto-tuning a Convolutional Network for NVIDIA GPU</a></li>
<li class="toctree-l3"><a class="reference internal" href="tune_relay_x86.html">为x86CPU 自动调优一个卷积网络</a></li>
<li class="toctree-l3"><a class="reference internal" href="tune_relay_arm.html">Auto-tuning a Convolutional Network for ARM CPU</a></li>
<li class="toctree-l3"><a class="reference internal" href="tune_relay_mobile_gpu.html">Auto-tuning a Convolutional Network for Mobile GPU</a></li>
</ul>
</li>
<li class="toctree-l2"><a class="reference internal" href="../tune_with_autoscheduler/index.html">Use AutoScheduler for Template-Free Scheduling</a></li>
<li class="toctree-l2"><a class="reference internal" href="../work_with_microtvm/index.html">Work With microTVM</a></li>
<li class="toctree-l2"><a class="reference internal" href="../extend_tvm/index.html">Extend TVM</a></li>
<li class="toctree-l2"><a class="reference internal" href="../profile/index.html">Profile Models</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../errors.html">Handle TVM Errors</a></li>
<li class="toctree-l2"><a class="reference internal" href="../../faq.html">常见提问</a></li>
</ul>
</li>
</ul>
<p class="caption" role="heading"><span class="caption-text">开发者引导</span></p>
<ul>
<li class="toctree-l1"><a class="reference internal" href="../../dev/tutorial/index.html">Developer Tutorial</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../dev/how_to/how_to.html">开发者指南</a></li>
</ul>
<p class="caption" role="heading"><span class="caption-text">架构指南</span></p>
<ul>
<li class="toctree-l1"><a class="reference internal" href="../../arch/index.html">Design and Architecture</a></li>
</ul>
<p class="caption" role="heading"><span class="caption-text">主题引导</span></p>
<ul>
<li class="toctree-l1"><a class="reference internal" href="../../topic/microtvm/index.html">microTVM：裸机使用TVM</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../topic/vta/index.html">VTA: Versatile Tensor Accelerator</a></li>
</ul>
<p class="caption" role="heading"><span class="caption-text">参考指南</span></p>
<ul>
<li class="toctree-l1"><a class="reference internal" href="../../reference/langref/index.html">语言参考</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../reference/api/python/index.html">Python API</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../reference/api/links.html">Other APIs</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../reference/publications.html">Publications</a></li>
<li class="toctree-l1"><a class="reference internal" href="../../genindex.html">索引</a></li>
</ul>

            
          
        </div>
        
      </div>
    </nav>

    <section data-toggle="wy-nav-shift" class="wy-nav-content-wrap">
      
      <nav class="wy-nav-top" aria-label="top navigation" data-toggle="wy-nav-top">
        
            <div class="togglemenu">

            </div>
            <div class="nav-content">
              <!-- tvm -->
              Table of content
            </div>
        
      </nav>


      <div class="wy-nav-content">
        
        <div class="rst-content">
        

          




















<div role="navigation" aria-label="breadcrumbs navigation">

  <ul class="wy-breadcrumbs">
    
      <li><a href="../../index.html">Docs</a> <span class="br-arrow">></span></li>
        
          <li><a href="../index.html">How To Guides</a> <span class="br-arrow">></span></li>
        
          <li><a href="index.html">Auto-Tune with Templates and AutoTVM</a> <span class="br-arrow">></span></li>
        
      <li>Tuning High Performance Convolution on NVIDIA GPUs</li>
    
    
      <li class="wy-breadcrumbs-aside">
        
            
            <a href="../../_sources/how_to/tune_with_autotvm/tune_conv2d_cuda.rst.txt" rel="nofollow"> <img src="../../_static//img/source.svg" alt="viewsource"/></a>
          
        
      </li>
    
  </ul>

  
  <hr/>
</div>
          <div role="main" class="document" itemscope="itemscope" itemtype="http://schema.org/Article">
           <div itemprop="articleBody">
            
  <div class="sphx-glr-download-link-note admonition note">
<p class="admonition-title">注解</p>
<p>点击 <a class="reference internal" href="#sphx-glr-download-how-to-tune-with-autotvm-tune-conv2d-cuda-py"><span class="std std-ref">此处</span></a> 获取完整示例代码</p>
</div>
<div class="sphx-glr-example-title section" id="tuning-high-performance-convolution-on-nvidia-gpus">
<span id="sphx-glr-how-to-tune-with-autotvm-tune-conv2d-cuda-py"></span><h1>Tuning High Performance Convolution on NVIDIA GPUs<a class="headerlink" href="#tuning-high-performance-convolution-on-nvidia-gpus" title="永久链接至标题">¶</a></h1>
<p><strong>作者</strong>: <a class="reference external" href="https://github.com/merrymercy">Lianmin Zheng</a></p>
<p>This is an advanced tutorial for writing high performance tunable template for
NVIDIA GPU. By running auto-tuner on this template, we can outperform the
vendor provided library CuDNN in many cases.</p>
<p>Note that this tutorial will not run on Windows or recent versions of macOS. To
get it to run, you will need to wrap the body of this tutorial in a <code class="code docutils literal notranslate"><span class="pre">if</span>
<span class="pre">__name__</span> <span class="pre">==</span> <span class="pre">&quot;__main__&quot;:</span></code> block.</p>
<div class="section" id="install-dependencies">
<h2>安装依赖<a class="headerlink" href="#install-dependencies" title="永久链接至标题">¶</a></h2>
<p>To use autotvm package in tvm, we need to install some extra dependencies.
(change “3” to “2” if you use python2):</p>
<div class="highlight-bash notranslate"><div class="highlight"><pre><span></span>pip3 install --user psutil xgboost tornado cloudpickle
</pre></div>
</div>
<p>To make TVM run faster in tuning, it is recommended to use cython
as FFI of tvm. In the root directory of tvm, execute</p>
<div class="highlight-bash notranslate"><div class="highlight"><pre><span></span>pip3 install --user cython
sudo make cython3
</pre></div>
</div>
<p>Now return to python code. Import packages.</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="kn">import</span> <span class="nn">logging</span>
<span class="kn">import</span> <span class="nn">sys</span>
<span class="kn">import</span> <span class="nn">numpy</span> <span class="k">as</span> <span class="nn">np</span>

<span class="kn">import</span> <span class="nn">tvm</span>
<span class="kn">from</span> <span class="nn">tvm</span> <span class="k">import</span> <span class="n">te</span><span class="p">,</span> <span class="n">topi</span><span class="p">,</span> <span class="n">testing</span>
<span class="kn">from</span> <span class="nn">tvm.topi.testing</span> <span class="k">import</span> <span class="n">conv2d_nchw_python</span>
<span class="kn">import</span> <span class="nn">tvm.testing</span>

<span class="kn">from</span> <span class="nn">tvm</span> <span class="k">import</span> <span class="n">autotvm</span>
</pre></div>
</div>
</div>
<div class="section" id="step-1-define-the-search-space">
<h2>Step 1:  Define the search space<a class="headerlink" href="#step-1-define-the-search-space" title="永久链接至标题">¶</a></h2>
<p>There are plenty of useful schedule primitives in tvm. You can also find
some tutorials that describe them in more details, such as
(1). <a class="reference internal" href="../optimize_operators/opt_conv_cuda.html#opt-conv-gpu"><span class="std std-ref">如何在GPU上优化卷积</span></a>
(2). <a class="reference external" href="https://tvm.apache.org/2017/08/22/Optimize-Deep-Learning-GPU-Operators-with-TVM-A-Depthwise-Convolution-Example">Optimizing DepthwiseConv on NVIDIA GPU</a></p>
<p>However, their implementations are manually tuned for some special input
shapes. In this section, we build a large enough space to cover
the techniques used in these tutorials. Then we rely on the efficient auto-tuner
to search through this space and pick some good configurations.</p>
<p>If you are familiar with writing cuda schedule, you can find the following
template is very general. Actually this template can be easily modified
to tune other operators such as depthwise convolution and gemm.
In order to fully understand this template, you should be familiar with
the schedule primitives and auto tuning API. You can refer to the above
tutorials and <a class="reference internal" href="../../tutorial/autotvm_matmul_x86.html#tutorial-autotvm-matmul-x86"><span class="std std-ref">autotvm tutorial</span></a></p>
<p>It is worth noting that the search space for a conv2d operator
can be very large (at the level of 10^9 for some input shapes)</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="nd">@autotvm</span><span class="o">.</span><span class="n">template</span><span class="p">(</span><span class="s2">&quot;tutorial/conv2d_no_batching&quot;</span><span class="p">)</span>
<span class="k">def</span> <span class="nf">conv2d_no_batching</span><span class="p">(</span><span class="n">N</span><span class="p">,</span> <span class="n">H</span><span class="p">,</span> <span class="n">W</span><span class="p">,</span> <span class="n">CO</span><span class="p">,</span> <span class="n">CI</span><span class="p">,</span> <span class="n">KH</span><span class="p">,</span> <span class="n">KW</span><span class="p">,</span> <span class="n">stride</span><span class="p">,</span> <span class="n">padding</span><span class="p">):</span>
    <span class="k">assert</span> <span class="n">N</span> <span class="o">==</span> <span class="mi">1</span><span class="p">,</span> <span class="s2">&quot;Only consider batch_size = 1 in this template&quot;</span>

    <span class="n">data</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">placeholder</span><span class="p">((</span><span class="n">N</span><span class="p">,</span> <span class="n">CI</span><span class="p">,</span> <span class="n">H</span><span class="p">,</span> <span class="n">W</span><span class="p">),</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;data&quot;</span><span class="p">)</span>
    <span class="n">kernel</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">placeholder</span><span class="p">((</span><span class="n">CO</span><span class="p">,</span> <span class="n">CI</span><span class="p">,</span> <span class="n">KH</span><span class="p">,</span> <span class="n">KW</span><span class="p">),</span> <span class="n">name</span><span class="o">=</span><span class="s2">&quot;kernel&quot;</span><span class="p">)</span>
    <span class="n">conv</span> <span class="o">=</span> <span class="n">topi</span><span class="o">.</span><span class="n">nn</span><span class="o">.</span><span class="n">conv2d_nchw</span><span class="p">(</span><span class="n">data</span><span class="p">,</span> <span class="n">kernel</span><span class="p">,</span> <span class="n">stride</span><span class="p">,</span> <span class="n">padding</span><span class="p">,</span> <span class="n">dilation</span><span class="o">=</span><span class="mi">1</span><span class="p">,</span> <span class="n">out_dtype</span><span class="o">=</span><span class="s2">&quot;float32&quot;</span><span class="p">)</span>
    <span class="n">s</span> <span class="o">=</span> <span class="n">te</span><span class="o">.</span><span class="n">create_schedule</span><span class="p">([</span><span class="n">conv</span><span class="o">.</span><span class="n">op</span><span class="p">])</span>

    <span class="c1">##### space definition begin #####</span>
    <span class="n">n</span><span class="p">,</span> <span class="n">f</span><span class="p">,</span> <span class="n">y</span><span class="p">,</span> <span class="n">x</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">conv</span><span class="p">]</span><span class="o">.</span><span class="n">op</span><span class="o">.</span><span class="n">axis</span>
    <span class="n">rc</span><span class="p">,</span> <span class="n">ry</span><span class="p">,</span> <span class="n">rx</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">conv</span><span class="p">]</span><span class="o">.</span><span class="n">op</span><span class="o">.</span><span class="n">reduce_axis</span>

    <span class="n">cfg</span> <span class="o">=</span> <span class="n">autotvm</span><span class="o">.</span><span class="n">get_config</span><span class="p">()</span>
    <span class="n">cfg</span><span class="o">.</span><span class="n">define_split</span><span class="p">(</span><span class="s2">&quot;tile_f&quot;</span><span class="p">,</span> <span class="n">f</span><span class="p">,</span> <span class="n">num_outputs</span><span class="o">=</span><span class="mi">4</span><span class="p">)</span>
    <span class="n">cfg</span><span class="o">.</span><span class="n">define_split</span><span class="p">(</span><span class="s2">&quot;tile_y&quot;</span><span class="p">,</span> <span class="n">y</span><span class="p">,</span> <span class="n">num_outputs</span><span class="o">=</span><span class="mi">4</span><span class="p">)</span>
    <span class="n">cfg</span><span class="o">.</span><span class="n">define_split</span><span class="p">(</span><span class="s2">&quot;tile_x&quot;</span><span class="p">,</span> <span class="n">x</span><span class="p">,</span> <span class="n">num_outputs</span><span class="o">=</span><span class="mi">4</span><span class="p">)</span>
    <span class="n">cfg</span><span class="o">.</span><span class="n">define_split</span><span class="p">(</span><span class="s2">&quot;tile_rc&quot;</span><span class="p">,</span> <span class="n">rc</span><span class="p">,</span> <span class="n">num_outputs</span><span class="o">=</span><span class="mi">3</span><span class="p">)</span>
    <span class="n">cfg</span><span class="o">.</span><span class="n">define_split</span><span class="p">(</span><span class="s2">&quot;tile_ry&quot;</span><span class="p">,</span> <span class="n">ry</span><span class="p">,</span> <span class="n">num_outputs</span><span class="o">=</span><span class="mi">3</span><span class="p">)</span>
    <span class="n">cfg</span><span class="o">.</span><span class="n">define_split</span><span class="p">(</span><span class="s2">&quot;tile_rx&quot;</span><span class="p">,</span> <span class="n">rx</span><span class="p">,</span> <span class="n">num_outputs</span><span class="o">=</span><span class="mi">3</span><span class="p">)</span>
    <span class="n">cfg</span><span class="o">.</span><span class="n">define_knob</span><span class="p">(</span><span class="s2">&quot;auto_unroll_max_step&quot;</span><span class="p">,</span> <span class="p">[</span><span class="mi">0</span><span class="p">,</span> <span class="mi">512</span><span class="p">,</span> <span class="mi">1500</span><span class="p">])</span>
    <span class="n">cfg</span><span class="o">.</span><span class="n">define_knob</span><span class="p">(</span><span class="s2">&quot;unroll_explicit&quot;</span><span class="p">,</span> <span class="p">[</span><span class="mi">0</span><span class="p">,</span> <span class="mi">1</span><span class="p">])</span>
    <span class="c1">##### space definition end #####</span>

    <span class="c1"># inline padding</span>
    <span class="n">pad_data</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">conv</span><span class="p">]</span><span class="o">.</span><span class="n">op</span><span class="o">.</span><span class="n">input_tensors</span><span class="p">[</span><span class="mi">0</span><span class="p">]</span>
    <span class="n">s</span><span class="p">[</span><span class="n">pad_data</span><span class="p">]</span><span class="o">.</span><span class="n">compute_inline</span><span class="p">()</span>
    <span class="n">data</span><span class="p">,</span> <span class="n">raw_data</span> <span class="o">=</span> <span class="n">pad_data</span><span class="p">,</span> <span class="n">data</span>

    <span class="n">output</span> <span class="o">=</span> <span class="n">conv</span>
    <span class="n">OL</span> <span class="o">=</span> <span class="n">s</span><span class="o">.</span><span class="n">cache_write</span><span class="p">(</span><span class="n">conv</span><span class="p">,</span> <span class="s2">&quot;local&quot;</span><span class="p">)</span>

    <span class="c1"># create cache stage</span>
    <span class="n">AA</span> <span class="o">=</span> <span class="n">s</span><span class="o">.</span><span class="n">cache_read</span><span class="p">(</span><span class="n">data</span><span class="p">,</span> <span class="s2">&quot;shared&quot;</span><span class="p">,</span> <span class="p">[</span><span class="n">OL</span><span class="p">])</span>
    <span class="n">WW</span> <span class="o">=</span> <span class="n">s</span><span class="o">.</span><span class="n">cache_read</span><span class="p">(</span><span class="n">kernel</span><span class="p">,</span> <span class="s2">&quot;shared&quot;</span><span class="p">,</span> <span class="p">[</span><span class="n">OL</span><span class="p">])</span>
    <span class="n">AL</span> <span class="o">=</span> <span class="n">s</span><span class="o">.</span><span class="n">cache_read</span><span class="p">(</span><span class="n">AA</span><span class="p">,</span> <span class="s2">&quot;local&quot;</span><span class="p">,</span> <span class="p">[</span><span class="n">OL</span><span class="p">])</span>
    <span class="n">WL</span> <span class="o">=</span> <span class="n">s</span><span class="o">.</span><span class="n">cache_read</span><span class="p">(</span><span class="n">WW</span><span class="p">,</span> <span class="s2">&quot;local&quot;</span><span class="p">,</span> <span class="p">[</span><span class="n">OL</span><span class="p">])</span>

    <span class="c1"># tile and bind spatial axes</span>
    <span class="n">n</span><span class="p">,</span> <span class="n">f</span><span class="p">,</span> <span class="n">y</span><span class="p">,</span> <span class="n">x</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">output</span><span class="p">]</span><span class="o">.</span><span class="n">op</span><span class="o">.</span><span class="n">axis</span>
    <span class="n">bf</span><span class="p">,</span> <span class="n">vf</span><span class="p">,</span> <span class="n">tf</span><span class="p">,</span> <span class="n">fi</span> <span class="o">=</span> <span class="n">cfg</span><span class="p">[</span><span class="s2">&quot;tile_f&quot;</span><span class="p">]</span><span class="o">.</span><span class="n">apply</span><span class="p">(</span><span class="n">s</span><span class="p">,</span> <span class="n">output</span><span class="p">,</span> <span class="n">f</span><span class="p">)</span>
    <span class="n">by</span><span class="p">,</span> <span class="n">vy</span><span class="p">,</span> <span class="n">ty</span><span class="p">,</span> <span class="n">yi</span> <span class="o">=</span> <span class="n">cfg</span><span class="p">[</span><span class="s2">&quot;tile_y&quot;</span><span class="p">]</span><span class="o">.</span><span class="n">apply</span><span class="p">(</span><span class="n">s</span><span class="p">,</span> <span class="n">output</span><span class="p">,</span> <span class="n">y</span><span class="p">)</span>
    <span class="n">bx</span><span class="p">,</span> <span class="n">vx</span><span class="p">,</span> <span class="n">tx</span><span class="p">,</span> <span class="n">xi</span> <span class="o">=</span> <span class="n">cfg</span><span class="p">[</span><span class="s2">&quot;tile_x&quot;</span><span class="p">]</span><span class="o">.</span><span class="n">apply</span><span class="p">(</span><span class="n">s</span><span class="p">,</span> <span class="n">output</span><span class="p">,</span> <span class="n">x</span><span class="p">)</span>
    <span class="n">kernel_scope</span> <span class="o">=</span> <span class="n">n</span>  <span class="c1"># this is the scope to attach global config inside this kernel</span>

    <span class="n">s</span><span class="p">[</span><span class="n">output</span><span class="p">]</span><span class="o">.</span><span class="n">bind</span><span class="p">(</span><span class="n">bf</span><span class="p">,</span> <span class="n">te</span><span class="o">.</span><span class="n">thread_axis</span><span class="p">(</span><span class="s2">&quot;blockIdx.z&quot;</span><span class="p">))</span>
    <span class="n">s</span><span class="p">[</span><span class="n">output</span><span class="p">]</span><span class="o">.</span><span class="n">bind</span><span class="p">(</span><span class="n">by</span><span class="p">,</span> <span class="n">te</span><span class="o">.</span><span class="n">thread_axis</span><span class="p">(</span><span class="s2">&quot;blockIdx.y&quot;</span><span class="p">))</span>
    <span class="n">s</span><span class="p">[</span><span class="n">output</span><span class="p">]</span><span class="o">.</span><span class="n">bind</span><span class="p">(</span><span class="n">bx</span><span class="p">,</span> <span class="n">te</span><span class="o">.</span><span class="n">thread_axis</span><span class="p">(</span><span class="s2">&quot;blockIdx.x&quot;</span><span class="p">))</span>
    <span class="n">s</span><span class="p">[</span><span class="n">output</span><span class="p">]</span><span class="o">.</span><span class="n">bind</span><span class="p">(</span><span class="n">vf</span><span class="p">,</span> <span class="n">te</span><span class="o">.</span><span class="n">thread_axis</span><span class="p">(</span><span class="s2">&quot;vthread&quot;</span><span class="p">))</span>
    <span class="n">s</span><span class="p">[</span><span class="n">output</span><span class="p">]</span><span class="o">.</span><span class="n">bind</span><span class="p">(</span><span class="n">vy</span><span class="p">,</span> <span class="n">te</span><span class="o">.</span><span class="n">thread_axis</span><span class="p">(</span><span class="s2">&quot;vthread&quot;</span><span class="p">))</span>
    <span class="n">s</span><span class="p">[</span><span class="n">output</span><span class="p">]</span><span class="o">.</span><span class="n">bind</span><span class="p">(</span><span class="n">vx</span><span class="p">,</span> <span class="n">te</span><span class="o">.</span><span class="n">thread_axis</span><span class="p">(</span><span class="s2">&quot;vthread&quot;</span><span class="p">))</span>
    <span class="n">s</span><span class="p">[</span><span class="n">output</span><span class="p">]</span><span class="o">.</span><span class="n">bind</span><span class="p">(</span><span class="n">tf</span><span class="p">,</span> <span class="n">te</span><span class="o">.</span><span class="n">thread_axis</span><span class="p">(</span><span class="s2">&quot;threadIdx.z&quot;</span><span class="p">))</span>
    <span class="n">s</span><span class="p">[</span><span class="n">output</span><span class="p">]</span><span class="o">.</span><span class="n">bind</span><span class="p">(</span><span class="n">ty</span><span class="p">,</span> <span class="n">te</span><span class="o">.</span><span class="n">thread_axis</span><span class="p">(</span><span class="s2">&quot;threadIdx.y&quot;</span><span class="p">))</span>
    <span class="n">s</span><span class="p">[</span><span class="n">output</span><span class="p">]</span><span class="o">.</span><span class="n">bind</span><span class="p">(</span><span class="n">tx</span><span class="p">,</span> <span class="n">te</span><span class="o">.</span><span class="n">thread_axis</span><span class="p">(</span><span class="s2">&quot;threadIdx.x&quot;</span><span class="p">))</span>
    <span class="n">s</span><span class="p">[</span><span class="n">output</span><span class="p">]</span><span class="o">.</span><span class="n">reorder</span><span class="p">(</span><span class="n">n</span><span class="p">,</span> <span class="n">bf</span><span class="p">,</span> <span class="n">by</span><span class="p">,</span> <span class="n">bx</span><span class="p">,</span> <span class="n">vf</span><span class="p">,</span> <span class="n">vy</span><span class="p">,</span> <span class="n">vx</span><span class="p">,</span> <span class="n">tf</span><span class="p">,</span> <span class="n">ty</span><span class="p">,</span> <span class="n">tx</span><span class="p">,</span> <span class="n">fi</span><span class="p">,</span> <span class="n">yi</span><span class="p">,</span> <span class="n">xi</span><span class="p">)</span>
    <span class="n">s</span><span class="p">[</span><span class="n">OL</span><span class="p">]</span><span class="o">.</span><span class="n">compute_at</span><span class="p">(</span><span class="n">s</span><span class="p">[</span><span class="n">output</span><span class="p">],</span> <span class="n">tx</span><span class="p">)</span>

    <span class="c1"># tile reduction axes</span>
    <span class="n">n</span><span class="p">,</span> <span class="n">f</span><span class="p">,</span> <span class="n">y</span><span class="p">,</span> <span class="n">x</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">OL</span><span class="p">]</span><span class="o">.</span><span class="n">op</span><span class="o">.</span><span class="n">axis</span>
    <span class="n">rc</span><span class="p">,</span> <span class="n">ry</span><span class="p">,</span> <span class="n">rx</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">OL</span><span class="p">]</span><span class="o">.</span><span class="n">op</span><span class="o">.</span><span class="n">reduce_axis</span>
    <span class="n">rco</span><span class="p">,</span> <span class="n">rcm</span><span class="p">,</span> <span class="n">rci</span> <span class="o">=</span> <span class="n">cfg</span><span class="p">[</span><span class="s2">&quot;tile_rc&quot;</span><span class="p">]</span><span class="o">.</span><span class="n">apply</span><span class="p">(</span><span class="n">s</span><span class="p">,</span> <span class="n">OL</span><span class="p">,</span> <span class="n">rc</span><span class="p">)</span>
    <span class="n">ryo</span><span class="p">,</span> <span class="n">rym</span><span class="p">,</span> <span class="n">ryi</span> <span class="o">=</span> <span class="n">cfg</span><span class="p">[</span><span class="s2">&quot;tile_rx&quot;</span><span class="p">]</span><span class="o">.</span><span class="n">apply</span><span class="p">(</span><span class="n">s</span><span class="p">,</span> <span class="n">OL</span><span class="p">,</span> <span class="n">ry</span><span class="p">)</span>
    <span class="n">rxo</span><span class="p">,</span> <span class="n">rxm</span><span class="p">,</span> <span class="n">rxi</span> <span class="o">=</span> <span class="n">cfg</span><span class="p">[</span><span class="s2">&quot;tile_ry&quot;</span><span class="p">]</span><span class="o">.</span><span class="n">apply</span><span class="p">(</span><span class="n">s</span><span class="p">,</span> <span class="n">OL</span><span class="p">,</span> <span class="n">rx</span><span class="p">)</span>
    <span class="n">s</span><span class="p">[</span><span class="n">OL</span><span class="p">]</span><span class="o">.</span><span class="n">reorder</span><span class="p">(</span><span class="n">rco</span><span class="p">,</span> <span class="n">ryo</span><span class="p">,</span> <span class="n">rxo</span><span class="p">,</span> <span class="n">rcm</span><span class="p">,</span> <span class="n">rym</span><span class="p">,</span> <span class="n">rxm</span><span class="p">,</span> <span class="n">rci</span><span class="p">,</span> <span class="n">ryi</span><span class="p">,</span> <span class="n">rxi</span><span class="p">,</span> <span class="n">n</span><span class="p">,</span> <span class="n">f</span><span class="p">,</span> <span class="n">y</span><span class="p">,</span> <span class="n">x</span><span class="p">)</span>

    <span class="n">s</span><span class="p">[</span><span class="n">AA</span><span class="p">]</span><span class="o">.</span><span class="n">compute_at</span><span class="p">(</span><span class="n">s</span><span class="p">[</span><span class="n">OL</span><span class="p">],</span> <span class="n">rxo</span><span class="p">)</span>
    <span class="n">s</span><span class="p">[</span><span class="n">WW</span><span class="p">]</span><span class="o">.</span><span class="n">compute_at</span><span class="p">(</span><span class="n">s</span><span class="p">[</span><span class="n">OL</span><span class="p">],</span> <span class="n">rxo</span><span class="p">)</span>
    <span class="n">s</span><span class="p">[</span><span class="n">AL</span><span class="p">]</span><span class="o">.</span><span class="n">compute_at</span><span class="p">(</span><span class="n">s</span><span class="p">[</span><span class="n">OL</span><span class="p">],</span> <span class="n">rxm</span><span class="p">)</span>
    <span class="n">s</span><span class="p">[</span><span class="n">WL</span><span class="p">]</span><span class="o">.</span><span class="n">compute_at</span><span class="p">(</span><span class="n">s</span><span class="p">[</span><span class="n">OL</span><span class="p">],</span> <span class="n">rxm</span><span class="p">)</span>

    <span class="c1"># cooperative fetching</span>
    <span class="k">for</span> <span class="n">load</span> <span class="ow">in</span> <span class="p">[</span><span class="n">AA</span><span class="p">,</span> <span class="n">WW</span><span class="p">]:</span>
        <span class="n">n</span><span class="p">,</span> <span class="n">f</span><span class="p">,</span> <span class="n">y</span><span class="p">,</span> <span class="n">x</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">load</span><span class="p">]</span><span class="o">.</span><span class="n">op</span><span class="o">.</span><span class="n">axis</span>
        <span class="n">fused</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">load</span><span class="p">]</span><span class="o">.</span><span class="n">fuse</span><span class="p">(</span><span class="n">n</span><span class="p">,</span> <span class="n">f</span><span class="p">,</span> <span class="n">y</span><span class="p">,</span> <span class="n">x</span><span class="p">)</span>
        <span class="n">tz</span><span class="p">,</span> <span class="n">fused</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">load</span><span class="p">]</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="n">fused</span><span class="p">,</span> <span class="n">nparts</span><span class="o">=</span><span class="n">cfg</span><span class="p">[</span><span class="s2">&quot;tile_f&quot;</span><span class="p">]</span><span class="o">.</span><span class="n">size</span><span class="p">[</span><span class="mi">2</span><span class="p">])</span>
        <span class="n">ty</span><span class="p">,</span> <span class="n">fused</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">load</span><span class="p">]</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="n">fused</span><span class="p">,</span> <span class="n">nparts</span><span class="o">=</span><span class="n">cfg</span><span class="p">[</span><span class="s2">&quot;tile_y&quot;</span><span class="p">]</span><span class="o">.</span><span class="n">size</span><span class="p">[</span><span class="mi">2</span><span class="p">])</span>
        <span class="n">tx</span><span class="p">,</span> <span class="n">fused</span> <span class="o">=</span> <span class="n">s</span><span class="p">[</span><span class="n">load</span><span class="p">]</span><span class="o">.</span><span class="n">split</span><span class="p">(</span><span class="n">fused</span><span class="p">,</span> <span class="n">nparts</span><span class="o">=</span><span class="n">cfg</span><span class="p">[</span><span class="s2">&quot;tile_x&quot;</span><span class="p">]</span><span class="o">.</span><span class="n">size</span><span class="p">[</span><span class="mi">2</span><span class="p">])</span>
        <span class="n">s</span><span class="p">[</span><span class="n">load</span><span class="p">]</span><span class="o">.</span><span class="n">bind</span><span class="p">(</span><span class="n">tz</span><span class="p">,</span> <span class="n">te</span><span class="o">.</span><span class="n">thread_axis</span><span class="p">(</span><span class="s2">&quot;threadIdx.z&quot;</span><span class="p">))</span>
        <span class="n">s</span><span class="p">[</span><span class="n">load</span><span class="p">]</span><span class="o">.</span><span class="n">bind</span><span class="p">(</span><span class="n">ty</span><span class="p">,</span> <span class="n">te</span><span class="o">.</span><span class="n">thread_axis</span><span class="p">(</span><span class="s2">&quot;threadIdx.y&quot;</span><span class="p">))</span>
        <span class="n">s</span><span class="p">[</span><span class="n">load</span><span class="p">]</span><span class="o">.</span><span class="n">bind</span><span class="p">(</span><span class="n">tx</span><span class="p">,</span> <span class="n">te</span><span class="o">.</span><span class="n">thread_axis</span><span class="p">(</span><span class="s2">&quot;threadIdx.x&quot;</span><span class="p">))</span>

    <span class="c1"># tune unroll</span>
    <span class="n">s</span><span class="p">[</span><span class="n">output</span><span class="p">]</span><span class="o">.</span><span class="n">pragma</span><span class="p">(</span><span class="n">kernel_scope</span><span class="p">,</span> <span class="s2">&quot;auto_unroll_max_step&quot;</span><span class="p">,</span> <span class="n">cfg</span><span class="p">[</span><span class="s2">&quot;auto_unroll_max_step&quot;</span><span class="p">]</span><span class="o">.</span><span class="n">val</span><span class="p">)</span>
    <span class="n">s</span><span class="p">[</span><span class="n">output</span><span class="p">]</span><span class="o">.</span><span class="n">pragma</span><span class="p">(</span><span class="n">kernel_scope</span><span class="p">,</span> <span class="s2">&quot;unroll_explicit&quot;</span><span class="p">,</span> <span class="n">cfg</span><span class="p">[</span><span class="s2">&quot;unroll_explicit&quot;</span><span class="p">]</span><span class="o">.</span><span class="n">val</span><span class="p">)</span>

    <span class="k">return</span> <span class="n">s</span><span class="p">,</span> <span class="p">[</span><span class="n">raw_data</span><span class="p">,</span> <span class="n">kernel</span><span class="p">,</span> <span class="n">conv</span><span class="p">]</span>
</pre></div>
</div>
</div>
<div class="section" id="step-2-search-through-the-space">
<h2>Step 2:  Search through the space<a class="headerlink" href="#step-2-search-through-the-space" title="永久链接至标题">¶</a></h2>
<p>We pick the last layer on resnet as test case.
Since our space is very large, <code class="code docutils literal notranslate"><span class="pre">XGBoostTuner</span></code> is most suitable
for our case. Here we only do 20 trials for demonstration.
In practice, making 1000 trials usually can find some good kernels
for this template</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="c1"># logging config (for printing tuning log to screen)</span>
<span class="n">logging</span><span class="o">.</span><span class="n">getLogger</span><span class="p">(</span><span class="s2">&quot;autotvm&quot;</span><span class="p">)</span><span class="o">.</span><span class="n">setLevel</span><span class="p">(</span><span class="n">logging</span><span class="o">.</span><span class="n">DEBUG</span><span class="p">)</span>
<span class="n">logging</span><span class="o">.</span><span class="n">getLogger</span><span class="p">(</span><span class="s2">&quot;autotvm&quot;</span><span class="p">)</span><span class="o">.</span><span class="n">addHandler</span><span class="p">(</span><span class="n">logging</span><span class="o">.</span><span class="n">StreamHandler</span><span class="p">(</span><span class="n">sys</span><span class="o">.</span><span class="n">stdout</span><span class="p">))</span>

<span class="c1"># the last layer in resnet</span>
<span class="n">N</span><span class="p">,</span> <span class="n">H</span><span class="p">,</span> <span class="n">W</span><span class="p">,</span> <span class="n">CO</span><span class="p">,</span> <span class="n">CI</span><span class="p">,</span> <span class="n">KH</span><span class="p">,</span> <span class="n">KW</span><span class="p">,</span> <span class="n">strides</span><span class="p">,</span> <span class="n">padding</span> <span class="o">=</span> <span class="mi">1</span><span class="p">,</span> <span class="mi">7</span><span class="p">,</span> <span class="mi">7</span><span class="p">,</span> <span class="mi">512</span><span class="p">,</span> <span class="mi">512</span><span class="p">,</span> <span class="mi">3</span><span class="p">,</span> <span class="mi">3</span><span class="p">,</span> <span class="p">(</span><span class="mi">1</span><span class="p">,</span> <span class="mi">1</span><span class="p">),</span> <span class="p">(</span><span class="mi">1</span><span class="p">,</span> <span class="mi">1</span><span class="p">)</span>
<span class="n">task</span> <span class="o">=</span> <span class="n">autotvm</span><span class="o">.</span><span class="n">task</span><span class="o">.</span><span class="n">create</span><span class="p">(</span>
    <span class="s2">&quot;tutorial/conv2d_no_batching&quot;</span><span class="p">,</span> <span class="n">args</span><span class="o">=</span><span class="p">(</span><span class="n">N</span><span class="p">,</span> <span class="n">H</span><span class="p">,</span> <span class="n">W</span><span class="p">,</span> <span class="n">CO</span><span class="p">,</span> <span class="n">CI</span><span class="p">,</span> <span class="n">KH</span><span class="p">,</span> <span class="n">KW</span><span class="p">,</span> <span class="n">strides</span><span class="p">,</span> <span class="n">padding</span><span class="p">),</span> <span class="n">target</span><span class="o">=</span><span class="s2">&quot;cuda&quot;</span>
<span class="p">)</span>
<span class="nb">print</span><span class="p">(</span><span class="n">task</span><span class="o">.</span><span class="n">config_space</span><span class="p">)</span>

<span class="c1"># Use local gpu, measure 10 times for every config to reduce variance</span>
<span class="c1"># The timeout of compiling a program is 10 seconds, the timeout for running is 4 seconds</span>
<span class="n">measure_option</span> <span class="o">=</span> <span class="n">autotvm</span><span class="o">.</span><span class="n">measure_option</span><span class="p">(</span>
    <span class="n">builder</span><span class="o">=</span><span class="n">autotvm</span><span class="o">.</span><span class="n">LocalBuilder</span><span class="p">(),</span>
    <span class="n">runner</span><span class="o">=</span><span class="n">autotvm</span><span class="o">.</span><span class="n">LocalRunner</span><span class="p">(</span><span class="n">repeat</span><span class="o">=</span><span class="mi">3</span><span class="p">,</span> <span class="n">min_repeat_ms</span><span class="o">=</span><span class="mi">100</span><span class="p">,</span> <span class="n">timeout</span><span class="o">=</span><span class="mi">4</span><span class="p">),</span>
<span class="p">)</span>

<span class="c1"># Begin tuning, log records to file `conv2d.log`</span>
<span class="c1"># During tuning we will also try many invalid configs, so you are expected to</span>
<span class="c1"># see many error reports. As long as you can see non-zero GFLOPS, it is okay.</span>
<span class="n">tuner</span> <span class="o">=</span> <span class="n">autotvm</span><span class="o">.</span><span class="n">tuner</span><span class="o">.</span><span class="n">XGBTuner</span><span class="p">(</span><span class="n">task</span><span class="p">)</span>
<span class="n">tuner</span><span class="o">.</span><span class="n">tune</span><span class="p">(</span>
    <span class="n">n_trial</span><span class="o">=</span><span class="mi">20</span><span class="p">,</span>
    <span class="n">measure_option</span><span class="o">=</span><span class="n">measure_option</span><span class="p">,</span>
    <span class="n">callbacks</span><span class="o">=</span><span class="p">[</span><span class="n">autotvm</span><span class="o">.</span><span class="n">callback</span><span class="o">.</span><span class="n">log_to_file</span><span class="p">(</span><span class="s2">&quot;conv2d.log&quot;</span><span class="p">)],</span>
<span class="p">)</span>
</pre></div>
</div>
<p class="sphx-glr-script-out">输出:</p>
<div class="sphx-glr-script-out highlight-none notranslate"><div class="highlight"><pre><span></span>ConfigSpace (len=10454400, space_map=
   0 tile_f: Split(policy=factors, product=512, num_outputs=4) len=220
   1 tile_y: Split(policy=factors, product=7, num_outputs=4) len=4
   2 tile_x: Split(policy=factors, product=7, num_outputs=4) len=4
   3 tile_rc: Split(policy=factors, product=512, num_outputs=3) len=55
   4 tile_ry: Split(policy=factors, product=3, num_outputs=3) len=3
   5 tile_rx: Split(policy=factors, product=3, num_outputs=3) len=3
   6 auto_unroll_max_step: OtherOption([0, 512, 1500]) len=3
   7 unroll_explicit: OtherOption([0, 1]) len=2
)
Get devices for measurement successfully!
No: 1   GFLOPS: 0.00/0.00       result: MeasureResult(costs=(InstantiationError(&#39;Traceback (most recent call last):\n  9: TVMFuncCall\n  8: tvm::runtime::TypedPackedFunc&lt;tvm::IRModule (tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)&gt;::AssignTypedLambda&lt;tvm::{lambda(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)#5}&gt;(tvm::{lambda(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)#5}, std::__cxx11::basic_string&lt;char, std::char_traits&lt;char&gt;, std::allocator&lt;char&gt; &gt;)::{lambda(tvm::runtime::TVMArgs const&amp;, tvm::runtime::TVMRetValue*)#1}::operator()(tvm::runtime::TVMArgs const, tvm::runtime::TVMRetValue) const\n  7: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, std::__cxx11::basic_string&lt;char, std::char_traits&lt;char&gt;, std::allocator&lt;char&gt; &gt; const&amp;, std::unordered_map&lt;tvm::te::Tensor, tvm::tir::Buffer, std::hash&lt;tvm::te::Tensor&gt;, std::equal_to&lt;tvm::te::Tensor&gt;, std::allocator&lt;std::pair&lt;tvm::te::Tensor const, tvm::tir::Buffer&gt; &gt; &gt; const&amp;, bool)\n  6: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array&lt;tvm::transform::Pass, void&gt;)\n  5: tvm::transform::Pass::operator()(tvm::IRModule) const\n  4: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  3: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  2: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  1: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  0: std::_Function_handler&lt;void (tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*), TVMFuncCreateFromCFunc::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#2}&gt;::_M_invoke(std::_Any_data const&amp;, tvm::runtime::TVMArgs&amp;&amp;, tvm::runtime::TVMRetValue*&amp;&amp;)\n  File &quot;/tvm/python/tvm/_ffi/_ctypes/packed_func.py&quot;, line 81, in cfun\n    rv = local_pyfunc(*pyargs)\n  File &quot;/tvm/python/tvm/autotvm/measure/measure_methods.py&quot;, line 814, in verify_pass\n    raise InstantiationError(&quot;Skipped because of invalid gpu kernel&quot;)\ntvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel&#39;,),), error_no=MeasureErrorNo.INSTANTIATION_ERROR, all_cost=0.16219329833984375, timestamp=1634034675.0486848)    [(&#39;tile_f&#39;, [-1, 4, 4, 2]), (&#39;tile_y&#39;, [-1, 7, 1, 1]), (&#39;tile_x&#39;, [-1, 1, 7, 1]), (&#39;tile_rc&#39;, [-1, 128, 2]), (&#39;tile_ry&#39;, [-1, 1, 3]), (&#39;tile_rx&#39;, [-1, 3, 1]), (&#39;auto_unroll_max_step&#39;, 512), (&#39;unroll_explicit&#39;, 1)],None,7999494
No: 2   GFLOPS: 0.00/0.00       result: MeasureResult(costs=(InstantiationError(&#39;Traceback (most recent call last):\n  9: TVMFuncCall\n  8: tvm::runtime::TypedPackedFunc&lt;tvm::IRModule (tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)&gt;::AssignTypedLambda&lt;tvm::{lambda(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)#5}&gt;(tvm::{lambda(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)#5}, std::__cxx11::basic_string&lt;char, std::char_traits&lt;char&gt;, std::allocator&lt;char&gt; &gt;)::{lambda(tvm::runtime::TVMArgs const&amp;, tvm::runtime::TVMRetValue*)#1}::operator()(tvm::runtime::TVMArgs const, tvm::runtime::TVMRetValue) const\n  7: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, std::__cxx11::basic_string&lt;char, std::char_traits&lt;char&gt;, std::allocator&lt;char&gt; &gt; const&amp;, std::unordered_map&lt;tvm::te::Tensor, tvm::tir::Buffer, std::hash&lt;tvm::te::Tensor&gt;, std::equal_to&lt;tvm::te::Tensor&gt;, std::allocator&lt;std::pair&lt;tvm::te::Tensor const, tvm::tir::Buffer&gt; &gt; &gt; const&amp;, bool)\n  6: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array&lt;tvm::transform::Pass, void&gt;)\n  5: tvm::transform::Pass::operator()(tvm::IRModule) const\n  4: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  3: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  2: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  1: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  0: std::_Function_handler&lt;void (tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*), TVMFuncCreateFromCFunc::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#2}&gt;::_M_invoke(std::_Any_data const&amp;, tvm::runtime::TVMArgs&amp;&amp;, tvm::runtime::TVMRetValue*&amp;&amp;)\n  File &quot;/tvm/python/tvm/_ffi/_ctypes/packed_func.py&quot;, line 81, in cfun\n    rv = local_pyfunc(*pyargs)\n  File &quot;/tvm/python/tvm/autotvm/measure/measure_methods.py&quot;, line 814, in verify_pass\n    raise InstantiationError(&quot;Skipped because of invalid gpu kernel&quot;)\ntvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel&#39;,),), error_no=MeasureErrorNo.INSTANTIATION_ERROR, all_cost=0.18407750129699707, timestamp=1634034675.048707)     [(&#39;tile_f&#39;, [-1, 1, 8, 2]), (&#39;tile_y&#39;, [-1, 1, 7, 1]), (&#39;tile_x&#39;, [-1, 1, 7, 1]), (&#39;tile_rc&#39;, [-1, 1, 64]), (&#39;tile_ry&#39;, [-1, 1, 3]), (&#39;tile_rx&#39;, [-1, 1, 3]), (&#39;auto_unroll_max_step&#39;, 1500), (&#39;unroll_explicit&#39;, 0)],None,5194279
No: 3   GFLOPS: 0.00/0.00       result: MeasureResult(costs=(InstantiationError(&#39;Traceback (most recent call last):\n  9: TVMFuncCall\n  8: tvm::runtime::TypedPackedFunc&lt;tvm::IRModule (tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)&gt;::AssignTypedLambda&lt;tvm::{lambda(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)#5}&gt;(tvm::{lambda(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)#5}, std::__cxx11::basic_string&lt;char, std::char_traits&lt;char&gt;, std::allocator&lt;char&gt; &gt;)::{lambda(tvm::runtime::TVMArgs const&amp;, tvm::runtime::TVMRetValue*)#1}::operator()(tvm::runtime::TVMArgs const, tvm::runtime::TVMRetValue) const\n  7: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, std::__cxx11::basic_string&lt;char, std::char_traits&lt;char&gt;, std::allocator&lt;char&gt; &gt; const&amp;, std::unordered_map&lt;tvm::te::Tensor, tvm::tir::Buffer, std::hash&lt;tvm::te::Tensor&gt;, std::equal_to&lt;tvm::te::Tensor&gt;, std::allocator&lt;std::pair&lt;tvm::te::Tensor const, tvm::tir::Buffer&gt; &gt; &gt; const&amp;, bool)\n  6: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array&lt;tvm::transform::Pass, void&gt;)\n  5: tvm::transform::Pass::operator()(tvm::IRModule) const\n  4: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  3: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  2: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  1: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  0: std::_Function_handler&lt;void (tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*), TVMFuncCreateFromCFunc::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#2}&gt;::_M_invoke(std::_Any_data const&amp;, tvm::runtime::TVMArgs&amp;&amp;, tvm::runtime::TVMRetValue*&amp;&amp;)\n  File &quot;/tvm/python/tvm/_ffi/_ctypes/packed_func.py&quot;, line 81, in cfun\n    rv = local_pyfunc(*pyargs)\n  File &quot;/tvm/python/tvm/autotvm/measure/measure_methods.py&quot;, line 814, in verify_pass\n    raise InstantiationError(&quot;Skipped because of invalid gpu kernel&quot;)\ntvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel&#39;,),), error_no=MeasureErrorNo.INSTANTIATION_ERROR, all_cost=0.18375635147094727, timestamp=1634034675.0487235)    [(&#39;tile_f&#39;, [-1, 8, 32, 1]), (&#39;tile_y&#39;, [-1, 1, 1, 7]), (&#39;tile_x&#39;, [-1, 1, 7, 1]), (&#39;tile_rc&#39;, [-1, 2, 64]), (&#39;tile_ry&#39;, [-1, 3, 1]), (&#39;tile_rx&#39;, [-1, 1, 1]), (&#39;auto_unroll_max_step&#39;, 1500), (&#39;unroll_explicit&#39;, 1)],None,9069983
No: 4   GFLOPS: 0.00/0.00       result: MeasureResult(costs=(InstantiationError(&#39;Traceback (most recent call last):\n  9: TVMFuncCall\n  8: tvm::runtime::TypedPackedFunc&lt;tvm::IRModule (tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)&gt;::AssignTypedLambda&lt;tvm::{lambda(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)#5}&gt;(tvm::{lambda(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)#5}, std::__cxx11::basic_string&lt;char, std::char_traits&lt;char&gt;, std::allocator&lt;char&gt; &gt;)::{lambda(tvm::runtime::TVMArgs const&amp;, tvm::runtime::TVMRetValue*)#1}::operator()(tvm::runtime::TVMArgs const, tvm::runtime::TVMRetValue) const\n  7: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, std::__cxx11::basic_string&lt;char, std::char_traits&lt;char&gt;, std::allocator&lt;char&gt; &gt; const&amp;, std::unordered_map&lt;tvm::te::Tensor, tvm::tir::Buffer, std::hash&lt;tvm::te::Tensor&gt;, std::equal_to&lt;tvm::te::Tensor&gt;, std::allocator&lt;std::pair&lt;tvm::te::Tensor const, tvm::tir::Buffer&gt; &gt; &gt; const&amp;, bool)\n  6: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array&lt;tvm::transform::Pass, void&gt;)\n  5: tvm::transform::Pass::operator()(tvm::IRModule) const\n  4: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  3: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  2: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  1: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  0: std::_Function_handler&lt;void (tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*), TVMFuncCreateFromCFunc::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#2}&gt;::_M_invoke(std::_Any_data const&amp;, tvm::runtime::TVMArgs&amp;&amp;, tvm::runtime::TVMRetValue*&amp;&amp;)\n  File &quot;/tvm/python/tvm/_ffi/_ctypes/packed_func.py&quot;, line 81, in cfun\n    rv = local_pyfunc(*pyargs)\n  File &quot;/tvm/python/tvm/autotvm/measure/measure_methods.py&quot;, line 814, in verify_pass\n    raise InstantiationError(&quot;Skipped because of invalid gpu kernel&quot;)\ntvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel&#39;,),), error_no=MeasureErrorNo.INSTANTIATION_ERROR, all_cost=0.12436294555664062, timestamp=1634034675.048734)     [(&#39;tile_f&#39;, [-1, 16, 16, 1]), (&#39;tile_y&#39;, [-1, 7, 1, 1]), (&#39;tile_x&#39;, [-1, 7, 1, 1]), (&#39;tile_rc&#39;, [-1, 16, 32]), (&#39;tile_ry&#39;, [-1, 1, 1]), (&#39;tile_rx&#39;, [-1, 3, 1]), (&#39;auto_unroll_max_step&#39;, 0), (&#39;unroll_explicit&#39;, 0)],None,736818
No: 5   GFLOPS: 0.00/0.00       result: MeasureResult(costs=(InstantiationError(&#39;Traceback (most recent call last):\n  9: TVMFuncCall\n  8: tvm::runtime::TypedPackedFunc&lt;tvm::IRModule (tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)&gt;::AssignTypedLambda&lt;tvm::{lambda(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)#5}&gt;(tvm::{lambda(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)#5}, std::__cxx11::basic_string&lt;char, std::char_traits&lt;char&gt;, std::allocator&lt;char&gt; &gt;)::{lambda(tvm::runtime::TVMArgs const&amp;, tvm::runtime::TVMRetValue*)#1}::operator()(tvm::runtime::TVMArgs const, tvm::runtime::TVMRetValue) const\n  7: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, std::__cxx11::basic_string&lt;char, std::char_traits&lt;char&gt;, std::allocator&lt;char&gt; &gt; const&amp;, std::unordered_map&lt;tvm::te::Tensor, tvm::tir::Buffer, std::hash&lt;tvm::te::Tensor&gt;, std::equal_to&lt;tvm::te::Tensor&gt;, std::allocator&lt;std::pair&lt;tvm::te::Tensor const, tvm::tir::Buffer&gt; &gt; &gt; const&amp;, bool)\n  6: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array&lt;tvm::transform::Pass, void&gt;)\n  5: tvm::transform::Pass::operator()(tvm::IRModule) const\n  4: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  3: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  2: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  1: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  0: std::_Function_handler&lt;void (tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*), TVMFuncCreateFromCFunc::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#2}&gt;::_M_invoke(std::_Any_data const&amp;, tvm::runtime::TVMArgs&amp;&amp;, tvm::runtime::TVMRetValue*&amp;&amp;)\n  File &quot;/tvm/python/tvm/_ffi/_ctypes/packed_func.py&quot;, line 81, in cfun\n    rv = local_pyfunc(*pyargs)\n  File &quot;/tvm/python/tvm/autotvm/measure/measure_methods.py&quot;, line 814, in verify_pass\n    raise InstantiationError(&quot;Skipped because of invalid gpu kernel&quot;)\ntvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel&#39;,),), error_no=MeasureErrorNo.INSTANTIATION_ERROR, all_cost=0.12966513633728027, timestamp=1634034675.048745)     [(&#39;tile_f&#39;, [-1, 4, 4, 32]), (&#39;tile_y&#39;, [-1, 1, 1, 7]), (&#39;tile_x&#39;, [-1, 1, 7, 1]), (&#39;tile_rc&#39;, [-1, 1, 128]), (&#39;tile_ry&#39;, [-1, 1, 3]), (&#39;tile_rx&#39;, [-1, 3, 1]), (&#39;auto_unroll_max_step&#39;, 512), (&#39;unroll_explicit&#39;, 0)],None,2885496
No: 6   GFLOPS: 174.85/174.85   result: MeasureResult(costs=(0.0013239794736842106,), error_no=MeasureErrorNo.NO_ERROR, all_cost=3.1157889366149902, timestamp=1634034686.8987) [(&#39;tile_f&#39;, [-1, 1, 1, 1]), (&#39;tile_y&#39;, [-1, 1, 1, 1]), (&#39;tile_x&#39;, [-1, 1, 7, 1]), (&#39;tile_rc&#39;, [-1, 4, 4]), (&#39;tile_ry&#39;, [-1, 3, 1]), (&#39;tile_rx&#39;, [-1, 1, 1]), (&#39;auto_unroll_max_step&#39;, 1500), (&#39;unroll_explicit&#39;, 0)],None,3754080
No: 7   GFLOPS: 0.00/174.85     result: MeasureResult(costs=(InstantiationError(&#39;Traceback (most recent call last):\n  9: TVMFuncCall\n  8: tvm::runtime::TypedPackedFunc&lt;tvm::IRModule (tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)&gt;::AssignTypedLambda&lt;tvm::{lambda(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)#5}&gt;(tvm::{lambda(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)#5}, std::__cxx11::basic_string&lt;char, std::char_traits&lt;char&gt;, std::allocator&lt;char&gt; &gt;)::{lambda(tvm::runtime::TVMArgs const&amp;, tvm::runtime::TVMRetValue*)#1}::operator()(tvm::runtime::TVMArgs const, tvm::runtime::TVMRetValue) const\n  7: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, std::__cxx11::basic_string&lt;char, std::char_traits&lt;char&gt;, std::allocator&lt;char&gt; &gt; const&amp;, std::unordered_map&lt;tvm::te::Tensor, tvm::tir::Buffer, std::hash&lt;tvm::te::Tensor&gt;, std::equal_to&lt;tvm::te::Tensor&gt;, std::allocator&lt;std::pair&lt;tvm::te::Tensor const, tvm::tir::Buffer&gt; &gt; &gt; const&amp;, bool)\n  6: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array&lt;tvm::transform::Pass, void&gt;)\n  5: tvm::transform::Pass::operator()(tvm::IRModule) const\n  4: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  3: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  2: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  1: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  0: std::_Function_handler&lt;void (tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*), TVMFuncCreateFromCFunc::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#2}&gt;::_M_invoke(std::_Any_data const&amp;, tvm::runtime::TVMArgs&amp;&amp;, tvm::runtime::TVMRetValue*&amp;&amp;)\n  File &quot;/tvm/python/tvm/_ffi/_ctypes/packed_func.py&quot;, line 81, in cfun\n    rv = local_pyfunc(*pyargs)\n  File &quot;/tvm/python/tvm/autotvm/measure/measure_methods.py&quot;, line 814, in verify_pass\n    raise InstantiationError(&quot;Skipped because of invalid gpu kernel&quot;)\ntvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel&#39;,),), error_no=MeasureErrorNo.INSTANTIATION_ERROR, all_cost=0.10683250427246094, timestamp=1634034675.5742111)    [(&#39;tile_f&#39;, [-1, 1, 16, 32]), (&#39;tile_y&#39;, [-1, 1, 1, 1]), (&#39;tile_x&#39;, [-1, 1, 7, 1]), (&#39;tile_rc&#39;, [-1, 256, 1]), (&#39;tile_ry&#39;, [-1, 1, 3]), (&#39;tile_rx&#39;, [-1, 3, 1]), (&#39;auto_unroll_max_step&#39;, 0), (&#39;unroll_explicit&#39;, 1)],None,6225319
No: 8   GFLOPS: 0.00/174.85     result: MeasureResult(costs=(InstantiationError(&#39;Traceback (most recent call last):\n  9: TVMFuncCall\n  8: tvm::runtime::TypedPackedFunc&lt;tvm::IRModule (tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)&gt;::AssignTypedLambda&lt;tvm::{lambda(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)#5}&gt;(tvm::{lambda(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)#5}, std::__cxx11::basic_string&lt;char, std::char_traits&lt;char&gt;, std::allocator&lt;char&gt; &gt;)::{lambda(tvm::runtime::TVMArgs const&amp;, tvm::runtime::TVMRetValue*)#1}::operator()(tvm::runtime::TVMArgs const, tvm::runtime::TVMRetValue) const\n  7: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, std::__cxx11::basic_string&lt;char, std::char_traits&lt;char&gt;, std::allocator&lt;char&gt; &gt; const&amp;, std::unordered_map&lt;tvm::te::Tensor, tvm::tir::Buffer, std::hash&lt;tvm::te::Tensor&gt;, std::equal_to&lt;tvm::te::Tensor&gt;, std::allocator&lt;std::pair&lt;tvm::te::Tensor const, tvm::tir::Buffer&gt; &gt; &gt; const&amp;, bool)\n  6: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array&lt;tvm::transform::Pass, void&gt;)\n  5: tvm::transform::Pass::operator()(tvm::IRModule) const\n  4: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  3: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  2: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  1: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  0: std::_Function_handler&lt;void (tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*), TVMFuncCreateFromCFunc::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#2}&gt;::_M_invoke(std::_Any_data const&amp;, tvm::runtime::TVMArgs&amp;&amp;, tvm::runtime::TVMRetValue*&amp;&amp;)\n  File &quot;/tvm/python/tvm/_ffi/_ctypes/packed_func.py&quot;, line 81, in cfun\n    rv = local_pyfunc(*pyargs)\n  File &quot;/tvm/python/tvm/autotvm/measure/measure_methods.py&quot;, line 814, in verify_pass\n    raise InstantiationError(&quot;Skipped because of invalid gpu kernel&quot;)\ntvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel&#39;,),), error_no=MeasureErrorNo.INSTANTIATION_ERROR, all_cost=0.1449904441833496, timestamp=1634034675.5742278)     [(&#39;tile_f&#39;, [-1, 2, 1, 32]), (&#39;tile_y&#39;, [-1, 1, 1, 1]), (&#39;tile_x&#39;, [-1, 1, 1, 1]), (&#39;tile_rc&#39;, [-1, 8, 64]), (&#39;tile_ry&#39;, [-1, 3, 1]), (&#39;tile_rx&#39;, [-1, 3, 1]), (&#39;auto_unroll_max_step&#39;, 0), (&#39;unroll_explicit&#39;, 0)],None,943546
No: 9   GFLOPS: 0.00/174.85     result: MeasureResult(costs=(InstantiationError(&#39;Traceback (most recent call last):\n  9: TVMFuncCall\n  8: tvm::runtime::TypedPackedFunc&lt;tvm::IRModule (tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)&gt;::AssignTypedLambda&lt;tvm::{lambda(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)#5}&gt;(tvm::{lambda(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)#5}, std::__cxx11::basic_string&lt;char, std::char_traits&lt;char&gt;, std::allocator&lt;char&gt; &gt;)::{lambda(tvm::runtime::TVMArgs const&amp;, tvm::runtime::TVMRetValue*)#1}::operator()(tvm::runtime::TVMArgs const, tvm::runtime::TVMRetValue) const\n  7: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, std::__cxx11::basic_string&lt;char, std::char_traits&lt;char&gt;, std::allocator&lt;char&gt; &gt; const&amp;, std::unordered_map&lt;tvm::te::Tensor, tvm::tir::Buffer, std::hash&lt;tvm::te::Tensor&gt;, std::equal_to&lt;tvm::te::Tensor&gt;, std::allocator&lt;std::pair&lt;tvm::te::Tensor const, tvm::tir::Buffer&gt; &gt; &gt; const&amp;, bool)\n  6: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array&lt;tvm::transform::Pass, void&gt;)\n  5: tvm::transform::Pass::operator()(tvm::IRModule) const\n  4: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  3: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  2: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  1: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  0: std::_Function_handler&lt;void (tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*), TVMFuncCreateFromCFunc::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#2}&gt;::_M_invoke(std::_Any_data const&amp;, tvm::runtime::TVMArgs&amp;&amp;, tvm::runtime::TVMRetValue*&amp;&amp;)\n  File &quot;/tvm/python/tvm/_ffi/_ctypes/packed_func.py&quot;, line 81, in cfun\n    rv = local_pyfunc(*pyargs)\n  File &quot;/tvm/python/tvm/autotvm/measure/measure_methods.py&quot;, line 814, in verify_pass\n    raise InstantiationError(&quot;Skipped because of invalid gpu kernel&quot;)\ntvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel&#39;,),), error_no=MeasureErrorNo.INSTANTIATION_ERROR, all_cost=0.14846134185791016, timestamp=1634034675.574238)     [(&#39;tile_f&#39;, [-1, 4, 16, 4]), (&#39;tile_y&#39;, [-1, 1, 1, 7]), (&#39;tile_x&#39;, [-1, 1, 1, 7]), (&#39;tile_rc&#39;, [-1, 16, 32]), (&#39;tile_ry&#39;, [-1, 1, 3]), (&#39;tile_rx&#39;, [-1, 3, 1]), (&#39;auto_unroll_max_step&#39;, 512), (&#39;unroll_explicit&#39;, 0)],None,2868708
No: 10  GFLOPS: 0.00/174.85     result: MeasureResult(costs=(TimeoutError(),), error_no=MeasureErrorNo.BUILD_TIMEOUT, all_cost=10, timestamp=1634034683.2688901)        [(&#39;tile_f&#39;, [-1, 32, 2, 4]), (&#39;tile_y&#39;, [-1, 1, 7, 1]), (&#39;tile_x&#39;, [-1, 1, 1, 7]), (&#39;tile_rc&#39;, [-1, 4, 2]), (&#39;tile_ry&#39;, [-1, 1, 1]), (&#39;tile_rx&#39;, [-1, 1, 3]), (&#39;auto_unroll_max_step&#39;, 1500), (&#39;unroll_explicit&#39;, 0)],None,4691833
No: 11  GFLOPS: 0.00/174.85     result: MeasureResult(costs=(InstantiationError(&#39;Traceback (most recent call last):\n  9: TVMFuncCall\n  8: tvm::runtime::TypedPackedFunc&lt;tvm::IRModule (tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)&gt;::AssignTypedLambda&lt;tvm::{lambda(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)#5}&gt;(tvm::{lambda(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)#5}, std::__cxx11::basic_string&lt;char, std::char_traits&lt;char&gt;, std::allocator&lt;char&gt; &gt;)::{lambda(tvm::runtime::TVMArgs const&amp;, tvm::runtime::TVMRetValue*)#1}::operator()(tvm::runtime::TVMArgs const, tvm::runtime::TVMRetValue) const\n  7: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, std::__cxx11::basic_string&lt;char, std::char_traits&lt;char&gt;, std::allocator&lt;char&gt; &gt; const&amp;, std::unordered_map&lt;tvm::te::Tensor, tvm::tir::Buffer, std::hash&lt;tvm::te::Tensor&gt;, std::equal_to&lt;tvm::te::Tensor&gt;, std::allocator&lt;std::pair&lt;tvm::te::Tensor const, tvm::tir::Buffer&gt; &gt; &gt; const&amp;, bool)\n  6: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array&lt;tvm::transform::Pass, void&gt;)\n  5: tvm::transform::Pass::operator()(tvm::IRModule) const\n  4: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  3: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  2: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  1: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  0: std::_Function_handler&lt;void (tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*), TVMFuncCreateFromCFunc::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#2}&gt;::_M_invoke(std::_Any_data const&amp;, tvm::runtime::TVMArgs&amp;&amp;, tvm::runtime::TVMRetValue*&amp;&amp;)\n  File &quot;/tvm/python/tvm/_ffi/_ctypes/packed_func.py&quot;, line 81, in cfun\n    rv = local_pyfunc(*pyargs)\n  File &quot;/tvm/python/tvm/autotvm/measure/measure_methods.py&quot;, line 814, in verify_pass\n    raise InstantiationError(&quot;Skipped because of invalid gpu kernel&quot;)\ntvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel&#39;,),), error_no=MeasureErrorNo.INSTANTIATION_ERROR, all_cost=0.13720202445983887, timestamp=1634034683.2689223)    [(&#39;tile_f&#39;, [-1, 1, 2, 64]), (&#39;tile_y&#39;, [-1, 1, 1, 1]), (&#39;tile_x&#39;, [-1, 1, 1, 1]), (&#39;tile_rc&#39;, [-1, 4, 4]), (&#39;tile_ry&#39;, [-1, 1, 3]), (&#39;tile_rx&#39;, [-1, 3, 1]), (&#39;auto_unroll_max_step&#39;, 0), (&#39;unroll_explicit&#39;, 0)],None,1042124
No: 12  GFLOPS: 0.00/174.85     result: MeasureResult(costs=(InstantiationError(&#39;Traceback (most recent call last):\n  9: TVMFuncCall\n  8: tvm::runtime::TypedPackedFunc&lt;tvm::IRModule (tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)&gt;::AssignTypedLambda&lt;tvm::{lambda(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)#5}&gt;(tvm::{lambda(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)#5}, std::__cxx11::basic_string&lt;char, std::char_traits&lt;char&gt;, std::allocator&lt;char&gt; &gt;)::{lambda(tvm::runtime::TVMArgs const&amp;, tvm::runtime::TVMRetValue*)#1}::operator()(tvm::runtime::TVMArgs const, tvm::runtime::TVMRetValue) const\n  7: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, std::__cxx11::basic_string&lt;char, std::char_traits&lt;char&gt;, std::allocator&lt;char&gt; &gt; const&amp;, std::unordered_map&lt;tvm::te::Tensor, tvm::tir::Buffer, std::hash&lt;tvm::te::Tensor&gt;, std::equal_to&lt;tvm::te::Tensor&gt;, std::allocator&lt;std::pair&lt;tvm::te::Tensor const, tvm::tir::Buffer&gt; &gt; &gt; const&amp;, bool)\n  6: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array&lt;tvm::transform::Pass, void&gt;)\n  5: tvm::transform::Pass::operator()(tvm::IRModule) const\n  4: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  3: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  2: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  1: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  0: std::_Function_handler&lt;void (tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*), TVMFuncCreateFromCFunc::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#2}&gt;::_M_invoke(std::_Any_data const&amp;, tvm::runtime::TVMArgs&amp;&amp;, tvm::runtime::TVMRetValue*&amp;&amp;)\n  File &quot;/tvm/python/tvm/_ffi/_ctypes/packed_func.py&quot;, line 81, in cfun\n    rv = local_pyfunc(*pyargs)\n  File &quot;/tvm/python/tvm/autotvm/measure/measure_methods.py&quot;, line 814, in verify_pass\n    raise InstantiationError(&quot;Skipped because of invalid gpu kernel&quot;)\ntvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel&#39;,),), error_no=MeasureErrorNo.INSTANTIATION_ERROR, all_cost=0.06606721878051758, timestamp=1634034683.2689369)    [(&#39;tile_f&#39;, [-1, 32, 1, 4]), (&#39;tile_y&#39;, [-1, 1, 1, 7]), (&#39;tile_x&#39;, [-1, 1, 7, 1]), (&#39;tile_rc&#39;, [-1, 32, 16]), (&#39;tile_ry&#39;, [-1, 1, 1]), (&#39;tile_rx&#39;, [-1, 1, 3]), (&#39;auto_unroll_max_step&#39;, 1500), (&#39;unroll_explicit&#39;, 1)],None,10013405
No: 13  GFLOPS: 0.00/174.85     result: MeasureResult(costs=(InstantiationError(&#39;Traceback (most recent call last):\n  9: TVMFuncCall\n  8: tvm::runtime::TypedPackedFunc&lt;tvm::IRModule (tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)&gt;::AssignTypedLambda&lt;tvm::{lambda(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)#5}&gt;(tvm::{lambda(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)#5}, std::__cxx11::basic_string&lt;char, std::char_traits&lt;char&gt;, std::allocator&lt;char&gt; &gt;)::{lambda(tvm::runtime::TVMArgs const&amp;, tvm::runtime::TVMRetValue*)#1}::operator()(tvm::runtime::TVMArgs const, tvm::runtime::TVMRetValue) const\n  7: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, std::__cxx11::basic_string&lt;char, std::char_traits&lt;char&gt;, std::allocator&lt;char&gt; &gt; const&amp;, std::unordered_map&lt;tvm::te::Tensor, tvm::tir::Buffer, std::hash&lt;tvm::te::Tensor&gt;, std::equal_to&lt;tvm::te::Tensor&gt;, std::allocator&lt;std::pair&lt;tvm::te::Tensor const, tvm::tir::Buffer&gt; &gt; &gt; const&amp;, bool)\n  6: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array&lt;tvm::transform::Pass, void&gt;)\n  5: tvm::transform::Pass::operator()(tvm::IRModule) const\n  4: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  3: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  2: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  1: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  0: std::_Function_handler&lt;void (tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*), TVMFuncCreateFromCFunc::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#2}&gt;::_M_invoke(std::_Any_data const&amp;, tvm::runtime::TVMArgs&amp;&amp;, tvm::runtime::TVMRetValue*&amp;&amp;)\n  File &quot;/tvm/python/tvm/_ffi/_ctypes/packed_func.py&quot;, line 81, in cfun\n    rv = local_pyfunc(*pyargs)\n  File &quot;/tvm/python/tvm/autotvm/measure/measure_methods.py&quot;, line 814, in verify_pass\n    raise InstantiationError(&quot;Skipped because of invalid gpu kernel&quot;)\ntvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel&#39;,),), error_no=MeasureErrorNo.INSTANTIATION_ERROR, all_cost=0.024266719818115234, timestamp=1634034683.2689455)   [(&#39;tile_f&#39;, [-1, 8, 8, 2]), (&#39;tile_y&#39;, [-1, 1, 1, 1]), (&#39;tile_x&#39;, [-1, 1, 7, 1]), (&#39;tile_rc&#39;, [-1, 4, 32]), (&#39;tile_ry&#39;, [-1, 3, 1]), (&#39;tile_rx&#39;, [-1, 1, 3]), (&#39;auto_unroll_max_step&#39;, 0), (&#39;unroll_explicit&#39;, 1)],None,6732082
No: 14  GFLOPS: 0.00/174.85     result: MeasureResult(costs=(InstantiationError(&#39;Traceback (most recent call last):\n  9: TVMFuncCall\n  8: tvm::runtime::TypedPackedFunc&lt;tvm::IRModule (tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)&gt;::AssignTypedLambda&lt;tvm::{lambda(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)#5}&gt;(tvm::{lambda(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)#5}, std::__cxx11::basic_string&lt;char, std::char_traits&lt;char&gt;, std::allocator&lt;char&gt; &gt;)::{lambda(tvm::runtime::TVMArgs const&amp;, tvm::runtime::TVMRetValue*)#1}::operator()(tvm::runtime::TVMArgs const, tvm::runtime::TVMRetValue) const\n  7: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, std::__cxx11::basic_string&lt;char, std::char_traits&lt;char&gt;, std::allocator&lt;char&gt; &gt; const&amp;, std::unordered_map&lt;tvm::te::Tensor, tvm::tir::Buffer, std::hash&lt;tvm::te::Tensor&gt;, std::equal_to&lt;tvm::te::Tensor&gt;, std::allocator&lt;std::pair&lt;tvm::te::Tensor const, tvm::tir::Buffer&gt; &gt; &gt; const&amp;, bool)\n  6: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array&lt;tvm::transform::Pass, void&gt;)\n  5: tvm::transform::Pass::operator()(tvm::IRModule) const\n  4: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  3: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  2: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  1: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  0: std::_Function_handler&lt;void (tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*), TVMFuncCreateFromCFunc::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#2}&gt;::_M_invoke(std::_Any_data const&amp;, tvm::runtime::TVMArgs&amp;&amp;, tvm::runtime::TVMRetValue*&amp;&amp;)\n  File &quot;/tvm/python/tvm/_ffi/_ctypes/packed_func.py&quot;, line 81, in cfun\n    rv = local_pyfunc(*pyargs)\n  File &quot;/tvm/python/tvm/autotvm/measure/measure_methods.py&quot;, line 814, in verify_pass\n    raise InstantiationError(&quot;Skipped because of invalid gpu kernel&quot;)\ntvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel&#39;,),), error_no=MeasureErrorNo.INSTANTIATION_ERROR, all_cost=0.14156842231750488, timestamp=1634034683.2689521)    [(&#39;tile_f&#39;, [-1, 2, 4, 32]), (&#39;tile_y&#39;, [-1, 7, 1, 1]), (&#39;tile_x&#39;, [-1, 1, 1, 1]), (&#39;tile_rc&#39;, [-1, 4, 128]), (&#39;tile_ry&#39;, [-1, 1, 3]), (&#39;tile_rx&#39;, [-1, 1, 1]), (&#39;auto_unroll_max_step&#39;, 512), (&#39;unroll_explicit&#39;, 1)],None,7536735
No: 15  GFLOPS: 0.00/174.85     result: MeasureResult(costs=(InstantiationError(&#39;Traceback (most recent call last):\n  9: TVMFuncCall\n  8: tvm::runtime::TypedPackedFunc&lt;tvm::IRModule (tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)&gt;::AssignTypedLambda&lt;tvm::{lambda(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)#5}&gt;(tvm::{lambda(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)#5}, std::__cxx11::basic_string&lt;char, std::char_traits&lt;char&gt;, std::allocator&lt;char&gt; &gt;)::{lambda(tvm::runtime::TVMArgs const&amp;, tvm::runtime::TVMRetValue*)#1}::operator()(tvm::runtime::TVMArgs const, tvm::runtime::TVMRetValue) const\n  7: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, std::__cxx11::basic_string&lt;char, std::char_traits&lt;char&gt;, std::allocator&lt;char&gt; &gt; const&amp;, std::unordered_map&lt;tvm::te::Tensor, tvm::tir::Buffer, std::hash&lt;tvm::te::Tensor&gt;, std::equal_to&lt;tvm::te::Tensor&gt;, std::allocator&lt;std::pair&lt;tvm::te::Tensor const, tvm::tir::Buffer&gt; &gt; &gt; const&amp;, bool)\n  6: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array&lt;tvm::transform::Pass, void&gt;)\n  5: tvm::transform::Pass::operator()(tvm::IRModule) const\n  4: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  3: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  2: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  1: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  0: std::_Function_handler&lt;void (tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*), TVMFuncCreateFromCFunc::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#2}&gt;::_M_invoke(std::_Any_data const&amp;, tvm::runtime::TVMArgs&amp;&amp;, tvm::runtime::TVMRetValue*&amp;&amp;)\n  File &quot;/tvm/python/tvm/_ffi/_ctypes/packed_func.py&quot;, line 81, in cfun\n    rv = local_pyfunc(*pyargs)\n  File &quot;/tvm/python/tvm/autotvm/measure/measure_methods.py&quot;, line 814, in verify_pass\n    raise InstantiationError(&quot;Skipped because of invalid gpu kernel&quot;)\ntvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel&#39;,),), error_no=MeasureErrorNo.INSTANTIATION_ERROR, all_cost=0.13557815551757812, timestamp=1634034683.2689598)    [(&#39;tile_f&#39;, [-1, 2, 1, 4]), (&#39;tile_y&#39;, [-1, 1, 1, 7]), (&#39;tile_x&#39;, [-1, 1, 1, 7]), (&#39;tile_rc&#39;, [-1, 128, 4]), (&#39;tile_ry&#39;, [-1, 1, 3]), (&#39;tile_rx&#39;, [-1, 1, 1]), (&#39;auto_unroll_max_step&#39;, 0), (&#39;unroll_explicit&#39;, 0)],None,482121
No: 16  GFLOPS: 0.00/174.85     result: MeasureResult(costs=(InstantiationError(&#39;Traceback (most recent call last):\n  9: TVMFuncCall\n  8: tvm::runtime::TypedPackedFunc&lt;tvm::IRModule (tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)&gt;::AssignTypedLambda&lt;tvm::{lambda(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)#5}&gt;(tvm::{lambda(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)#5}, std::__cxx11::basic_string&lt;char, std::char_traits&lt;char&gt;, std::allocator&lt;char&gt; &gt;)::{lambda(tvm::runtime::TVMArgs const&amp;, tvm::runtime::TVMRetValue*)#1}::operator()(tvm::runtime::TVMArgs const, tvm::runtime::TVMRetValue) const\n  7: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, std::__cxx11::basic_string&lt;char, std::char_traits&lt;char&gt;, std::allocator&lt;char&gt; &gt; const&amp;, std::unordered_map&lt;tvm::te::Tensor, tvm::tir::Buffer, std::hash&lt;tvm::te::Tensor&gt;, std::equal_to&lt;tvm::te::Tensor&gt;, std::allocator&lt;std::pair&lt;tvm::te::Tensor const, tvm::tir::Buffer&gt; &gt; &gt; const&amp;, bool)\n  6: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array&lt;tvm::transform::Pass, void&gt;)\n  5: tvm::transform::Pass::operator()(tvm::IRModule) const\n  4: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  3: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  2: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  1: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  0: std::_Function_handler&lt;void (tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*), TVMFuncCreateFromCFunc::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#2}&gt;::_M_invoke(std::_Any_data const&amp;, tvm::runtime::TVMArgs&amp;&amp;, tvm::runtime::TVMRetValue*&amp;&amp;)\n  File &quot;/tvm/python/tvm/_ffi/_ctypes/packed_func.py&quot;, line 81, in cfun\n    rv = local_pyfunc(*pyargs)\n  File &quot;/tvm/python/tvm/autotvm/measure/measure_methods.py&quot;, line 814, in verify_pass\n    raise InstantiationError(&quot;Skipped because of invalid gpu kernel&quot;)\ntvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel&#39;,),), error_no=MeasureErrorNo.INSTANTIATION_ERROR, all_cost=0.05985879898071289, timestamp=1634034683.2689662)    [(&#39;tile_f&#39;, [-1, 2, 1, 16]), (&#39;tile_y&#39;, [-1, 1, 7, 1]), (&#39;tile_x&#39;, [-1, 7, 1, 1]), (&#39;tile_rc&#39;, [-1, 32, 8]), (&#39;tile_ry&#39;, [-1, 1, 3]), (&#39;tile_rx&#39;, [-1, 3, 1]), (&#39;auto_unroll_max_step&#39;, 512), (&#39;unroll_explicit&#39;, 0)],None,2824525
No: 17  GFLOPS: 0.00/174.85     result: MeasureResult(costs=(InstantiationError(&#39;Traceback (most recent call last):\n  9: TVMFuncCall\n  8: tvm::runtime::TypedPackedFunc&lt;tvm::IRModule (tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)&gt;::AssignTypedLambda&lt;tvm::{lambda(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)#5}&gt;(tvm::{lambda(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)#5}, std::__cxx11::basic_string&lt;char, std::char_traits&lt;char&gt;, std::allocator&lt;char&gt; &gt;)::{lambda(tvm::runtime::TVMArgs const&amp;, tvm::runtime::TVMRetValue*)#1}::operator()(tvm::runtime::TVMArgs const, tvm::runtime::TVMRetValue) const\n  7: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, std::__cxx11::basic_string&lt;char, std::char_traits&lt;char&gt;, std::allocator&lt;char&gt; &gt; const&amp;, std::unordered_map&lt;tvm::te::Tensor, tvm::tir::Buffer, std::hash&lt;tvm::te::Tensor&gt;, std::equal_to&lt;tvm::te::Tensor&gt;, std::allocator&lt;std::pair&lt;tvm::te::Tensor const, tvm::tir::Buffer&gt; &gt; &gt; const&amp;, bool)\n  6: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array&lt;tvm::transform::Pass, void&gt;)\n  5: tvm::transform::Pass::operator()(tvm::IRModule) const\n  4: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  3: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  2: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  1: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  0: std::_Function_handler&lt;void (tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*), TVMFuncCreateFromCFunc::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#2}&gt;::_M_invoke(std::_Any_data const&amp;, tvm::runtime::TVMArgs&amp;&amp;, tvm::runtime::TVMRetValue*&amp;&amp;)\n  File &quot;/tvm/python/tvm/_ffi/_ctypes/packed_func.py&quot;, line 81, in cfun\n    rv = local_pyfunc(*pyargs)\n  File &quot;/tvm/python/tvm/autotvm/measure/measure_methods.py&quot;, line 814, in verify_pass\n    raise InstantiationError(&quot;Skipped because of invalid gpu kernel&quot;)\ntvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel&#39;,),), error_no=MeasureErrorNo.INSTANTIATION_ERROR, all_cost=0.01901412010192871, timestamp=1634034683.268973)     [(&#39;tile_f&#39;, [-1, 64, 1, 1]), (&#39;tile_y&#39;, [-1, 1, 1, 1]), (&#39;tile_x&#39;, [-1, 7, 1, 1]), (&#39;tile_rc&#39;, [-1, 8, 8]), (&#39;tile_ry&#39;, [-1, 1, 3]), (&#39;tile_rx&#39;, [-1, 3, 1]), (&#39;auto_unroll_max_step&#39;, 1500), (&#39;unroll_explicit&#39;, 0)],None,4559286
No: 18  GFLOPS: 0.00/174.85     result: MeasureResult(costs=(InstantiationError(&#39;Traceback (most recent call last):\n  9: TVMFuncCall\n  8: tvm::runtime::TypedPackedFunc&lt;tvm::IRModule (tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)&gt;::AssignTypedLambda&lt;tvm::{lambda(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)#5}&gt;(tvm::{lambda(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, tvm::runtime::String const&amp;, tvm::runtime::Map&lt;tvm::te::Tensor, tvm::tir::Buffer, void, void&gt; const&amp;, bool)#5}, std::__cxx11::basic_string&lt;char, std::char_traits&lt;char&gt;, std::allocator&lt;char&gt; &gt;)::{lambda(tvm::runtime::TVMArgs const&amp;, tvm::runtime::TVMRetValue*)#1}::operator()(tvm::runtime::TVMArgs const, tvm::runtime::TVMRetValue) const\n  7: tvm::LowerSchedule(tvm::te::Schedule, tvm::runtime::Array&lt;tvm::runtime::ObjectRef, void&gt; const&amp;, std::__cxx11::basic_string&lt;char, std::char_traits&lt;char&gt;, std::allocator&lt;char&gt; &gt; const&amp;, std::unordered_map&lt;tvm::te::Tensor, tvm::tir::Buffer, std::hash&lt;tvm::te::Tensor&gt;, std::equal_to&lt;tvm::te::Tensor&gt;, std::allocator&lt;std::pair&lt;tvm::te::Tensor const, tvm::tir::Buffer&gt; &gt; &gt; const&amp;, bool)\n  6: tvm::LowerWithPassList(tvm::IRModule, tvm::runtime::Array&lt;tvm::transform::Pass, void&gt;)\n  5: tvm::transform::Pass::operator()(tvm::IRModule) const\n  4: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  3: tvm::transform::SequentialNode::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  2: tvm::transform::Pass::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  1: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&amp;) const\n  0: std::_Function_handler&lt;void (tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*), TVMFuncCreateFromCFunc::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#2}&gt;::_M_invoke(std::_Any_data const&amp;, tvm::runtime::TVMArgs&amp;&amp;, tvm::runtime::TVMRetValue*&amp;&amp;)\n  File &quot;/tvm/python/tvm/_ffi/_ctypes/packed_func.py&quot;, line 81, in cfun\n    rv = local_pyfunc(*pyargs)\n  File &quot;/tvm/python/tvm/autotvm/measure/measure_methods.py&quot;, line 814, in verify_pass\n    raise InstantiationError(&quot;Skipped because of invalid gpu kernel&quot;)\ntvm.autotvm.task.space.InstantiationError: Skipped because of invalid gpu kernel&#39;,),), error_no=MeasureErrorNo.INSTANTIATION_ERROR, all_cost=0.054665565490722656, timestamp=1634034683.2689805)   [(&#39;tile_f&#39;, [-1, 1, 32, 16]), (&#39;tile_y&#39;, [-1, 1, 1, 1]), (&#39;tile_x&#39;, [-1, 7, 1, 1]), (&#39;tile_rc&#39;, [-1, 1, 512]), (&#39;tile_ry&#39;, [-1, 3, 1]), (&#39;tile_rx&#39;, [-1, 3, 1]), (&#39;auto_unroll_max_step&#39;, 1500), (&#39;unroll_explicit&#39;, 1)],None,9677544
No: 19  GFLOPS: 0.86/174.85     result: MeasureResult(costs=(0.268915582,), error_no=MeasureErrorNo.NO_ERROR, all_cost=6.550044059753418, timestamp=1634034699.3013427) [(&#39;tile_f&#39;, [-1, 8, 2, 16]), (&#39;tile_y&#39;, [-1, 7, 1, 1]), (&#39;tile_x&#39;, [-1, 7, 1, 1]), (&#39;tile_rc&#39;, [-1, 1, 1]), (&#39;tile_ry&#39;, [-1, 1, 1]), (&#39;tile_rx&#39;, [-1, 1, 3]), (&#39;auto_unroll_max_step&#39;, 0), (&#39;unroll_explicit&#39;, 1)],None,6390073
No: 20  GFLOPS: 299.79/299.79   result: MeasureResult(costs=(0.0007722233221153847,), error_no=MeasureErrorNo.NO_ERROR, all_cost=1.7415635585784912, timestamp=1634034700.2218585)      [(&#39;tile_f&#39;, [-1, 1, 4, 1]), (&#39;tile_y&#39;, [-1, 1, 1, 1]), (&#39;tile_x&#39;, [-1, 7, 1, 1]), (&#39;tile_rc&#39;, [-1, 4, 1]), (&#39;tile_ry&#39;, [-1, 1, 1]), (&#39;tile_rx&#39;, [-1, 1, 3]), (&#39;auto_unroll_max_step&#39;, 1500), (&#39;unroll_explicit&#39;, 1)],None,9881539
</pre></div>
</div>
<p>Finally we can inspect the best config from log file, check correctness,
and measure running time.</p>
<div class="highlight-default notranslate"><div class="highlight"><pre><span></span><span class="c1"># inspect the best config</span>
<span class="n">dispatch_context</span> <span class="o">=</span> <span class="n">autotvm</span><span class="o">.</span><span class="n">apply_history_best</span><span class="p">(</span><span class="s2">&quot;conv2d.log&quot;</span><span class="p">)</span>
<span class="n">best_config</span> <span class="o">=</span> <span class="n">dispatch_context</span><span class="o">.</span><span class="n">query</span><span class="p">(</span><span class="n">task</span><span class="o">.</span><span class="n">target</span><span class="p">,</span> <span class="n">task</span><span class="o">.</span><span class="n">workload</span><span class="p">)</span>
<span class="nb">print</span><span class="p">(</span><span class="s2">&quot;</span><span class="se">\n</span><span class="s2">Best config:&quot;</span><span class="p">)</span>
<span class="nb">print</span><span class="p">(</span><span class="n">best_config</span><span class="p">)</span>

<span class="c1"># apply history best from log file</span>
<span class="k">with</span> <span class="n">autotvm</span><span class="o">.</span><span class="n">apply_history_best</span><span class="p">(</span><span class="s2">&quot;conv2d.log&quot;</span><span class="p">):</span>
    <span class="k">with</span> <span class="n">tvm</span><span class="o">.</span><span class="n">target</span><span class="o">.</span><span class="n">Target</span><span class="p">(</span><span class="s2">&quot;cuda&quot;</span><span class="p">):</span>
        <span class="n">s</span><span class="p">,</span> <span class="n">arg_bufs</span> <span class="o">=</span> <span class="n">conv2d_no_batching</span><span class="p">(</span><span class="n">N</span><span class="p">,</span> <span class="n">H</span><span class="p">,</span> <span class="n">W</span><span class="p">,</span> <span class="n">CO</span><span class="p">,</span> <span class="n">CI</span><span class="p">,</span> <span class="n">KH</span><span class="p">,</span> <span class="n">KW</span><span class="p">,</span> <span class="n">strides</span><span class="p">,</span> <span class="n">padding</span><span class="p">)</span>
        <span class="n">func</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">build</span><span class="p">(</span><span class="n">s</span><span class="p">,</span> <span class="n">arg_bufs</span><span class="p">)</span>

<span class="c1"># check correctness</span>
<span class="n">a_np</span> <span class="o">=</span> <span class="n">np</span><span class="o">.</span><span class="n">random</span><span class="o">.</span><span class="n">uniform</span><span class="p">(</span><span class="n">size</span><span class="o">=</span><span class="p">(</span><span class="n">N</span><span class="p">,</span> <span class="n">CI</span><span class="p">,</span> <span class="n">H</span><span class="p">,</span> <span class="n">W</span><span class="p">))</span><span class="o">.</span><span class="n">astype</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">float32</span><span class="p">)</span>
<span class="n">w_np</span> <span class="o">=</span> <span class="n">np</span><span class="o">.</span><span class="n">random</span><span class="o">.</span><span class="n">uniform</span><span class="p">(</span><span class="n">size</span><span class="o">=</span><span class="p">(</span><span class="n">CO</span><span class="p">,</span> <span class="n">CI</span><span class="p">,</span> <span class="n">KH</span><span class="p">,</span> <span class="n">KW</span><span class="p">))</span><span class="o">.</span><span class="n">astype</span><span class="p">(</span><span class="n">np</span><span class="o">.</span><span class="n">float32</span><span class="p">)</span>
<span class="n">c_np</span> <span class="o">=</span> <span class="n">conv2d_nchw_python</span><span class="p">(</span><span class="n">a_np</span><span class="p">,</span> <span class="n">w_np</span><span class="p">,</span> <span class="n">strides</span><span class="p">,</span> <span class="n">padding</span><span class="p">)</span>

<span class="n">dev</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">cuda</span><span class="p">()</span>
<span class="n">a_tvm</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">nd</span><span class="o">.</span><span class="n">array</span><span class="p">(</span><span class="n">a_np</span><span class="p">,</span> <span class="n">device</span><span class="o">=</span><span class="n">dev</span><span class="p">)</span>
<span class="n">w_tvm</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">nd</span><span class="o">.</span><span class="n">array</span><span class="p">(</span><span class="n">w_np</span><span class="p">,</span> <span class="n">device</span><span class="o">=</span><span class="n">dev</span><span class="p">)</span>
<span class="n">c_tvm</span> <span class="o">=</span> <span class="n">tvm</span><span class="o">.</span><span class="n">nd</span><span class="o">.</span><span class="n">empty</span><span class="p">(</span><span class="n">c_np</span><span class="o">.</span><span class="n">shape</span><span class="p">,</span> <span class="n">device</span><span class="o">=</span><span class="n">dev</span><span class="p">)</span>
<span class="n">func</span><span class="p">(</span><span class="n">a_tvm</span><span class="p">,</span> <span class="n">w_tvm</span><span class="p">,</span> <span class="n">c_tvm</span><span class="p">)</span>

<span class="n">tvm</span><span class="o">.</span><span class="n">testing</span><span class="o">.</span><span class="n">assert_allclose</span><span class="p">(</span><span class="n">c_np</span><span class="p">,</span> <span class="n">c_tvm</span><span class="o">.</span><span class="n">numpy</span><span class="p">(),</span> <span class="n">rtol</span><span class="o">=</span><span class="mf">1e-2</span><span class="p">)</span>

<span class="c1"># Evaluate running time. Here we choose a large repeat number (400) to reduce the noise</span>
<span class="c1"># and the overhead of kernel launch. You can also use nvprof to validate the result.</span>
<span class="n">evaluator</span> <span class="o">=</span> <span class="n">func</span><span class="o">.</span><span class="n">time_evaluator</span><span class="p">(</span><span class="n">func</span><span class="o">.</span><span class="n">entry_name</span><span class="p">,</span> <span class="n">dev</span><span class="p">,</span> <span class="n">number</span><span class="o">=</span><span class="mi">400</span><span class="p">)</span>
<span class="nb">print</span><span class="p">(</span><span class="s2">&quot;Time cost of this operator: </span><span class="si">%f</span><span class="s2">&quot;</span> <span class="o">%</span> <span class="n">evaluator</span><span class="p">(</span><span class="n">a_tvm</span><span class="p">,</span> <span class="n">w_tvm</span><span class="p">,</span> <span class="n">c_tvm</span><span class="p">)</span><span class="o">.</span><span class="n">mean</span><span class="p">)</span>
</pre></div>
</div>
<p class="sphx-glr-script-out">输出:</p>
<div class="sphx-glr-script-out highlight-none notranslate"><div class="highlight"><pre><span></span>Best config:
[(&#39;tile_f&#39;, [-1, 1, 4, 1]), (&#39;tile_y&#39;, [-1, 1, 1, 1]), (&#39;tile_x&#39;, [-1, 7, 1, 1]), (&#39;tile_rc&#39;, [-1, 4, 1]), (&#39;tile_ry&#39;, [-1, 1, 1]), (&#39;tile_rx&#39;, [-1, 1, 3]), (&#39;auto_unroll_max_step&#39;, 1500), (&#39;unroll_explicit&#39;, 1)],None,9881539
Time cost of this operator: 0.000938
</pre></div>
</div>
<div class="sphx-glr-footer class sphx-glr-footer-example docutils container" id="sphx-glr-download-how-to-tune-with-autotvm-tune-conv2d-cuda-py">
<div class="sphx-glr-download docutils container">
<p><a class="reference download internal" download="" href="../../_downloads/6ad550da5092845382b1197f58a93816/tune_conv2d_cuda.py"><code class="xref download docutils literal notranslate"><span class="pre">Python</span> <span class="pre">源码下载:</span> <span class="pre">tune_conv2d_cuda.py</span></code></a></p>
</div>
<div class="sphx-glr-download docutils container">
<p><a class="reference download internal" download="" href="../../_downloads/732ed130cbc15432e737da8cc47e1734/tune_conv2d_cuda.ipynb"><code class="xref download docutils literal notranslate"><span class="pre">Jupyter</span> <span class="pre">notebook</span> <span class="pre">下载:</span> <span class="pre">tune_conv2d_cuda.ipynb</span></code></a></p>
</div>
</div>
<p class="sphx-glr-signature"><a class="reference external" href="https://sphinx-gallery.github.io">Gallery generated by Sphinx-Gallery</a></p>
</div>
</div>


           </div>
           
          </div>
          

<footer>

    <div class="rst-footer-buttons" role="navigation" aria-label="footer navigation">
      
        <a href="tune_relay_cuda.html" class="btn btn-neutral float-right" title="Auto-tuning a Convolutional Network for NVIDIA GPU" accesskey="n" rel="next">下一个 <span class="fa fa-arrow-circle-right"></span></a>
      
      
        <a href="index.html" class="btn btn-neutral float-left" title="Auto-Tune with Templates and AutoTVM" accesskey="p" rel="prev"><span class="fa fa-arrow-circle-left"></span> 上一个</a>
      
    </div>

<div id="button" class="backtop"><img src="../../_static//img/right.svg" alt="backtop"/> </div>
<section class="footerSec">
    <div class="footerHeader">
      <ul class="d-flex align-md-items-center justify-content-between flex-column flex-md-row">
        <li class="copywrite d-flex align-items-center">
          <h5 id="copy-right-info">© 2020 Apache Software Foundation | All right reserved</h5>
        </li>
      </ul>

    </div>

    <ul>
      <li class="footernote">Copyright © 2020 The Apache Software Foundation. Apache TVM, Apache, the Apache feather, and the Apache TVM project logo are either trademarks or registered trademarks of the Apache Software Foundation.</li>
    </ul>

</section>
</footer>
        </div>
      </div>

    </section>

  </div>
  

    <script src="https://cdnjs.cloudflare.com/ajax/libs/popper.js/1.12.9/umd/popper.min.js" integrity="sha384-ApNbgh9B+Y1QKtv3Rn7W3mgPxhU9K/ScQsAP7hUibX39j7fakFPskvXusvfa0b4Q" crossorigin="anonymous"></script>
    <script src="https://maxcdn.bootstrapcdn.com/bootstrap/4.0.0/js/bootstrap.min.js" integrity="sha384-JZR6Spejh4U02d8jOt6vLEHfe/JQGiRRSQQxSfFWpi1MquVdAyjUar5+76PVCmYl" crossorigin="anonymous"></script>

  </body>
  <script type="text/javascript">
      jQuery(function () {
          SphinxRtdTheme.Navigation.enable(true);
      });
  </script>

  
  
    
    <!-- Theme Analytics -->
    <script>
    (function(i,s,o,g,r,a,m){i['GoogleAnalyticsObject']=r;i[r]=i[r]||function(){
      (i[r].q=i[r].q||[]).push(arguments)},i[r].l=1*new Date();a=s.createElement(o),
      m=s.getElementsByTagName(o)[0];a.async=1;a.src=g;m.parentNode.insertBefore(a,m)
    })(window,document,'script','https://www.google-analytics.com/analytics.js','ga');

    ga('create', 'UA-75982049-2', 'auto');
    ga('send', 'pageview');
    </script>

    
   

</body>
</html>